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

Change-Id: I35862740e33def19512bc82122da3a6c31f47a85

bin/hipcc.bat

@@ -1,2 +1,2 @@
-@IF DEFINED HIP_PATH (set HIPCC="%HIP_PATH%/bin/hipcc") ELSE (set HIPCC="%~dp0/hipcc")
+set HIPCC="%~dp0/hipcc"
 @perl %HIPCC% %*

bin/hipconfig.bat

@@ -1,2 +1,2 @@
-@IF DEFINED HIP_PATH (set HIPCONFIG="%HIP_PATH%/bin/hipconfig") ELSE (set HIPCONFIG="%~dp0/hipconfig")
+set HIPCONFIG="%~dp0/hipconfig"
 @perl %HIPCONFIG% %*

cmake/FindHIP.cmake

@@ -50,8 +50,8 @@ mark_as_advanced(HIP_HOST_COMPILATION_CPP)
 
 get_filename_component(_IMPORT_PREFIX "${CMAKE_CURRENT_LIST_DIR}/../" REALPATH)
 
-# HIP is supported on Linux only
-if(UNIX AND NOT APPLE AND NOT CYGWIN)
+# HIP is currently not supported for apple
+if(NOT APPLE)
     # Search for HIP installation
     if(NOT HIP_ROOT_DIR)
         # Search in user specified path first
@@ -94,7 +94,6 @@ if(UNIX AND NOT APPLE AND NOT CYGWIN)
         # Now search in default paths
         find_program(HIP_HIPCC_EXECUTABLE hipcc)
     endif()
-    mark_as_advanced(HIP_HIPCC_EXECUTABLE)
 
     # Find HIPCONFIG executable
     find_program(
@@ -113,7 +112,12 @@ if(UNIX AND NOT APPLE AND NOT CYGWIN)
         # Now search in default paths
         find_program(HIP_HIPCONFIG_EXECUTABLE hipconfig)
     endif()
+    if(NOT UNIX)
+        set(HIP_HIPCONFIG_EXECUTABLE "${HIP_HIPCONFIG_EXECUTABLE}.bat")
+        set(HIP_HIPCC_EXECUTABLE "${HIP_HIPCC_EXECUTABLE}.bat")
+    endif()
     mark_as_advanced(HIP_HIPCONFIG_EXECUTABLE)
+    mark_as_advanced(HIP_HIPCC_EXECUTABLE)
 
     # Find HIPCC_CMAKE_LINKER_HELPER executable
     find_program(

include/hip/hip_bfloat16.h

@@ -56,7 +56,7 @@ struct hip_bfloat16
 
     enum truncate_t
     {
-        truncate
+        truncate_0
     };
 
     __host__ __device__ hip_bfloat16() = default;

include/hip/hip_runtime_api.h

@@ -445,6 +445,8 @@ typedef enum hipDeviceAttribute_t {
                                                                 ///< hipStreamWaitValue64(), '0' otherwise.
     hipDeviceAttributeImageSupport,                             ///< '1' if Device supports image, '0' otherwise.
 
+    hipDeviceAttributeMultiprocessorBoostCount,                 ///< All available boost compute units for the device
+
     hipDeviceAttributeAmdSpecificEnd = 19999,
     hipDeviceAttributeVendorSpecificBegin = 20000,
     // Extended attributes for vendors
@@ -4302,6 +4304,11 @@ typedef enum hipStreamUpdateCaptureDependenciesFlags {
   hipStreamSetCaptureDependencies,      ///< Replace the dependency set with the new nodes
 } hipStreamUpdateCaptureDependenciesFlags;
 
+typedef enum hipGraphInstantiateFlags {
+  hipGraphInstantiateFlagAutoFreeOnLaunch =
+      1,  ///< Automatically free memory allocated in a graph before relaunching.
+} hipGraphInstantiateFlags;
+
 /**
  * @brief Begins graph capture on a stream.
  *
@@ -4399,6 +4406,31 @@ hipError_t hipStreamUpdateCaptureDependencies(hipStream_t stream, hipGraphNode_t
                                               size_t numDependencies,
                                               unsigned int flags __dparm(0));
 
+/**
+ * @brief Enqueues a host function call in a stream.
+ *
+ * @param [in] stream - stream to enqueue work to.
+ * @param [in] fn - function to call once operations enqueued preceeding are complete.
+ * @param [in] userData - User-specified data to be passed to the function.
+ * @returns #hipSuccess, #hipErrorInvalidResourceHandle, #hipErrorInvalidValue,
+ * #hipErrorNotSupported
+ * @warning : This API is marked as beta, meaning, while this is feature complete,
+ * it is still open to changes and may have outstanding issues.
+ */
+hipError_t hipLaunchHostFunc(hipStream_t stream, hipHostFn_t fn, void* userData);
+
+/**
+ * @brief Swaps the stream capture mode of a thread.
+ *
+ * @param [in] mode - Pointer to mode value to swap with the current mode
+ * @returns #hipSuccess, #hipErrorInvalidValue
+ *
+ * @warning : This API is marked as beta, meaning, while this is feature complete,
+ * it is still open to changes and may have outstanding issues.
+ *
+ */
+hipError_t hipThreadExchangeStreamCaptureMode(hipStreamCaptureMode* mode);
+
 /**
  * @brief Creates a graph
  *

samples/0_Intro/module_api/launchKernelHcc.cpp

@@ -69,9 +69,6 @@ int main() {
     HIP_CHECK(hipModuleLoad(&Module, fileName));
     HIP_CHECK(hipModuleGetFunction(&Function, Module, kernel_name));
 
-    uint32_t len = LEN;
-    uint32_t one = 1;
-
     struct {
         void* _Ad;
         void* _Bd;

samples/1_Utils/hipBusBandwidth/ResultDatabase.cpp

@@ -401,7 +401,6 @@ void ResultDatabase::DumpCsv(string fileName) {
 // ****************************************************************************
 
 bool ResultDatabase::IsFileEmpty(string fileName) {
-    bool fileEmpty;
 
     ifstream file(fileName.c_str());
 
@@ -409,6 +408,7 @@ bool ResultDatabase::IsFileEmpty(string fileName) {
     if (!file.good()) {
         return true;
     } else {
+        bool fileEmpty;
         fileEmpty = (bool)(file.peek() == ifstream::traits_type::eof());
         file.close();
 

samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp

@@ -252,9 +252,9 @@ void RunBenchmark_H2D(ResultDatabase& resultDB) {
 
         case MallocUnpinned:
             if (p_alignedhost) {
-                delete[] hostMem;
-            } else {
                 free(hostMem);
+            } else {
+                delete[] hostMem;
             }
             break;
 

samples/1_Utils/hipCommander/ResultDatabase.cpp

@@ -393,7 +393,6 @@ void ResultDatabase::DumpCsv(string fileName) {
 // ****************************************************************************
 
 bool ResultDatabase::IsFileEmpty(string fileName) {
-    bool fileEmpty;
 
     ifstream file(fileName.c_str());
 
@@ -401,6 +400,7 @@ bool ResultDatabase::IsFileEmpty(string fileName) {
     if (!file.good()) {
         return true;
     } else {
+        bool fileEmpty;
         fileEmpty = (bool)(file.peek() == ifstream::traits_type::eof());
         file.close();
 

samples/1_Utils/hipCommander/hipCommander.cpp

@@ -282,7 +282,7 @@ class Command {
 // HCC optimizes away fully NULL kernel calls, so run one that is nearly null:
 class ModuleKernelCommand : public Command {
    public:
-    ModuleKernelCommand(CommandStream* cmdStream, const std::vector<std::string> args)
+    ModuleKernelCommand(CommandStream* cmdStream, const std::vector<std::string>& args)
         : Command(cmdStream, args), _stream(cmdStream->currentStream()) {
         hipModule_t module;
         HIPCHECK(hipModuleLoad(&module, FILENAME));
@@ -316,7 +316,7 @@ class ModuleKernelCommand : public Command {
 class KernelCommand : public Command {
    public:
     enum Type { Null, VectorAdd };
-    KernelCommand(CommandStream* cmdStream, const std::vector<std::string> args, Type kind)
+    KernelCommand(CommandStream* cmdStream, const std::vector<std::string>& args, Type kind)
         : Command(cmdStream, args), _kind(kind), _stream(cmdStream->currentStream()){};
     ~KernelCommand(){};
 
@@ -390,7 +390,7 @@ class CopyCommand : public Command {
     };
 
 
-    void dealloc(void* p, MemType memType) {
+    static void dealloc(void* p, MemType memType) {
         if (memType == Device) {
             HIPCHECK(hipFree(p));
         } else if (memType == PinnedHost) {
@@ -433,7 +433,7 @@ class StreamSyncCommand : public Command {
     StreamSyncCommand(CommandStream* cmdStream, const std::vector<std::string>& args)
         : Command(cmdStream, args), _stream(cmdStream->currentStream()){};
 
-    const char* help() { return "synchronizes the current stream"; };
+    static const char* help() { return "synchronizes the current stream"; };
 
 
     void run() override { HIPCHECK(hipStreamSynchronize(_stream)); };
@@ -537,8 +537,8 @@ CopyCommand::CopyCommand(CommandStream* cmdStream, const std::vector<std::string
                          hipMemcpyKind kind, bool isAsync, bool isPinnedHost)
     : Command(cmdStream, args),
       _isAsync(isAsync),
-      _kind(kind),
-      _stream(cmdStream->currentStream()) {
+      _stream(cmdStream->currentStream()),
+      _kind(kind) {
     switch (kind) {
         case hipMemcpyDeviceToHost:
             _srcType = Device;

samples/2_Cookbook/14_gpu_arch/gpuarch.cpp

@@ -48,17 +48,14 @@ __global__ void incrementKernel(int32_t* in, int32_t* out, int32_t value, size_t
 
 int main() {
   int32_t incrementValue = 10;
-  // Host pointers
-  int32_t* hInput = nullptr;
-  int32_t* hOutput = nullptr;
   // Device pointers
   int32_t* dInput = nullptr;
   int32_t* dOutput = nullptr;
 
   size_t NBytes = SIZE * sizeof(int32_t);
-
-  hInput = static_cast<int32_t*>(malloc(NBytes));
-  hOutput = static_cast<int32_t*>(malloc(NBytes));
+  // Host pointers
+  int32_t* hInput = static_cast<int32_t*>(malloc(NBytes));
+  int32_t* hOutput = static_cast<int32_t*>(malloc(NBytes));
 
   HIP_STATUS_CHECK(hipMalloc(&dInput, NBytes));
   HIP_STATUS_CHECK(hipMalloc(&dOutput, NBytes));
@@ -95,4 +92,4 @@ int main() {
     std::cout << "success\n";
   }
   return 0;
-}
+}

samples/2_Cookbook/15_static_library/device_functions/hipMain2.cpp

@@ -58,8 +58,8 @@ void run_test2() {
 
     HIP_ASSERT(hipFree(A_d));
     HIP_ASSERT(hipFree(B_d));
-    free(A_h);
-    free(B_h);
+    delete [] A_h;
+    delete [] B_h;
     std::cout << "Test Passed!\n";
 }
 

tests/catch/stress/memory/hipMallocManagedStress.cc

@@ -1,5 +1,5 @@
 /*
-   Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved.
+   Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
    Permission is hereby granted, free of charge, to any person obtaining a copy
    of this software and associated documentation files (the "Software"), to deal
    in the Software without restriction, including without limitation the rights
@@ -19,6 +19,15 @@
 
 //  The following test case allocation, host access, device access of HMM
 //   memory from size 1 to 10KB
+/*  Test Case Description:
+    1) Testing allocation, host access, device access of HMM
+     memory from size 1 to 10KB
+    2) The following test case tests the behavior of kernel with a HMM memory
+        and hipMalloc memory
+    3) The following test case tests when the same Hmm memory is used for
+       launching multiple different kernels will results in any issue
+    4) Testing the allocation of/scenarios around max possible memory
+     */
 
 #include <hip_test_common.hh>
 #include <hip_test_checkers.hh>
@@ -35,6 +44,86 @@ __global__ void KrnlWth2MemTypesC(unsigned char *Hmm, unsigned char *Dptr,
 }
 static bool IfTestPassed = true;
 
+// Kernel functions
+__global__ void KrnlWth2MemTypes(int *Hmm, int *Dptr, size_t n) {
+  size_t index = blockIdx.x * blockDim.x + threadIdx.x;
+  for (size_t i = index; i < n; i++) {
+    Hmm[i] = Dptr[i] + 10;
+  }
+}
+
+__global__ void KernelMulAdd_MngdMem(int *Hmm, size_t n) {
+  size_t index = blockIdx.x * blockDim.x + threadIdx.x;
+  size_t stride = blockDim.x * gridDim.x;
+  for (size_t i = index; i < n; i += stride) {
+    Hmm[i] = Hmm[i] * 2 + 10;
+  }
+}
+
+__global__ void KernelMul_MngdMem(int *Hmm, int *Dptr, size_t n) {
+  size_t index = blockIdx.x * blockDim.x + threadIdx.x;
+  size_t stride = blockDim.x * gridDim.x;
+  for (size_t i = index; i < n; i += stride) {
+    Hmm[i] = Dptr[i] * 10;
+  }
+}
+static bool IfTestPassed = true;
+
+static void LaunchKrnl4(size_t NumElms, int InitVal) {
+  int *Hmm = NULL, *Dptr = NULL, blockSize = 64, DataMismatch = 0;
+  hipStream_t strm;
+  HIP_CHECK(hipStreamCreate(&strm));
+  HIP_CHECK(hipMallocManaged(&Hmm, (sizeof(int) * NumElms)));
+  HIP_CHECK(hipMalloc(&Dptr, (sizeof(int) * NumElms)));
+  int *Hstptr = reinterpret_cast<int*>(new int[NumElms]);
+  for (size_t i = 0; i < NumElms; ++i) {
+    Hstptr[i] = InitVal;
+  }
+  HIP_CHECK(hipMemcpy(Dptr, Hstptr, (NumElms * sizeof(int)),
+                      hipMemcpyHostToDevice));
+  dim3 dimBlock(blockSize, 1, 1);
+  dim3 dimGrid((NumElms + blockSize -1)/blockSize, 1, 1);
+  KrnlWth2MemTypes<<<dimGrid, dimBlock, 0, strm>>>(Hmm, Dptr, NumElms);
+  HIP_CHECK(hipStreamSynchronize(strm));
+  for (size_t i = 0; i < NumElms; ++i) {
+    if (Hmm[i] != (InitVal + 10)) {
+      DataMismatch++;
+    }
+  }
+  if (DataMismatch != 0) {
+    INFO("Data Mismatch observed after the Kernel: KrnlWth2MemTypes!!\n");
+    REQUIRE(false);
+  }
+  DataMismatch = 0;
+  KernelMul_MngdMem<<<dimGrid, dimBlock, 0, strm>>>(Hmm, Dptr, NumElms);
+  HIP_CHECK(hipStreamSynchronize(strm));
+  // Verifying the result
+  for (size_t i = 0; i < NumElms; ++i) {
+    if (Hmm[i] != (InitVal * 10)) {
+      DataMismatch++;
+    }
+  }
+  if (DataMismatch != 0) {
+    INFO("Data Mismatch observedafter the Kernel: KernelMul_MngdMem!!\n");
+    REQUIRE(false);
+  }
+  DataMismatch = 0;
+  KernelMulAdd_MngdMem<<<dimGrid, dimBlock, 0, strm>>>(Hmm, NumElms);
+  HIP_CHECK(hipStreamSynchronize(strm));
+  // Verifying the result
+
+  for (size_t i = 0; i < NumElms; ++i) {
+    if (Hmm[i] != (InitVal * 10 * 2 + 10)) {
+      DataMismatch++;
+    }
+  }
+  if (DataMismatch != 0) {
+    INFO("Data Mismatch observedafter the Kernel: KernelMul_MngdMem!!\n");
+    REQUIRE(false);
+  }
+  delete[] Hstptr;
+}
+
 static int HmmAttrPrint() {
   int managed = 0;
   INFO("The following are the attribute values related to HMM for"
@@ -104,3 +193,139 @@ TEST_CASE("Unit_hipMallocManaged_MultiSize") {
   }
 }
 
+// The following test case tests the behavior of kernel with a HMM memory and
+// hipMalloc memory
+
+TEST_CASE("Unit_hipMallocManaged_KrnlWth2MemTypes") {
+  IfTestPassed = true;
+  int *Hmm = NULL, *Dptr = NULL, InitVal = 123;
+  size_t NumElms = (1024 * 1024);
+  int *Hptr = new int[NumElms], blockSize = 64, DataMismatch = 0;
+  int managed =  HmmAttrPrint();
+  if (managed == 1) {
+    hipStream_t strm;
+    HIP_CHECK(hipStreamCreate(&strm));
+    HIP_CHECK(hipMallocManaged(&Hmm, sizeof(int) * NumElms));
+    HIP_CHECK(hipMalloc(&Dptr, sizeof(int) * NumElms));
+    for (size_t i = 0; i < NumElms; ++i) {
+      Hmm[i] = 0;
+      Hptr[i] = InitVal;
+    }
+    HIP_CHECK(hipMemcpy(Dptr, Hptr, sizeof(int) * NumElms,
+                        hipMemcpyHostToDevice));
+    dim3 dimBlock(blockSize, 1, 1);
+    dim3 dimGrid((NumElms + blockSize -1)/blockSize, 1, 1);
+    KrnlWth2MemTypes<<<dimGrid, dimBlock, 0, strm>>>(Hmm, Dptr, NumElms);
+    HIP_CHECK(hipStreamSynchronize(strm));
+    // Verifying the results
+    for (size_t k = 0; k < NumElms; ++k) {
+      if (Hmm[k] != (InitVal + 10)) {
+        DataMismatch++;
+      }
+    }
+    if (DataMismatch != 0) {
+      WARN("DataMismatch observed!\n");
+      IfTestPassed = false;
+    }
+
+    HIP_CHECK(hipFree(Hmm));
+    HIP_CHECK(hipFree(Dptr));
+    delete[] Hptr;
+    REQUIRE(IfTestPassed);
+  } else {
+    SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
+           "attribute. Hence skipping the testing with Pass result.\n");
+  }
+}
+
+
+// The following test case tests when the same Hmm memory is used for
+// launching multiple different kernels will results in any issue
+TEST_CASE("Unit_hipMallocManaged_MultiKrnlHmmAccess") {
+  int managed = HmmAttrPrint();
+  if (managed) {
+    int InitVal = 123, NumElms = (1024 * 1024);
+    LaunchKrnl4(NumElms, InitVal);
+  } else {
+    SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
+           "attribute. Hence skipping the testing with Pass result.\n");
+  }
+}
+
+// Testing the allocation of/scenarios around max possible memory
+TEST_CASE("Unit_hipMallocManaged_ExtremeSizes") {
+  int managed = HmmAttrPrint();
+  if (managed == 1) {
+    bool IfTestPassed = true;
+    hipError_t err;
+    void *Hmm = NULL;
+    size_t totalDevMem = 0, freeDevMem = 0;
+    int NumDevs = 0;
+    HIP_CHECK(hipGetDeviceCount(&NumDevs));
+    // Testing allocation of extreme and unusual mem values
+    for (int i = 0; i < NumDevs; i++) {
+      HIP_CHECK(hipSetDevice(i));
+      HIP_CHECK(hipMemGetInfo(&freeDevMem, &totalDevMem));
+      err = hipMallocManaged(&Hmm, 1, hipMemAttachGlobal);
+      if (hipSuccess == err) {
+        HIP_CHECK(hipFree(Hmm));
+      } else {
+        WARN("Observed error while allocating memory on GPU: " << i);
+        WARN(" size 1 with");
+        WARN(" hipMallocManaged() api with flag 'hipMemAttachGlobal'\n");
+        WARN("Error: " << hipGetErrorString(err));
+        IfTestPassed = false;
+      }
+      err = hipMallocManaged(&Hmm, freeDevMem, hipMemAttachGlobal);
+      if (hipSuccess == err) {
+        HIP_CHECK(hipFree(Hmm));
+      } else {
+        WARN("Observed error while allocating max free memory on GPU: " << i);
+        WARN(" with hipMallocManaged() api with flag 'hipMemAttachGlobal'\n");
+        WARN("Error: " << hipGetErrorString(err));
+        IfTestPassed = false;
+      }
+      err = hipMallocManaged(&Hmm, (freeDevMem - 1), hipMemAttachGlobal);
+      if (hipSuccess == err) {
+        HIP_CHECK(hipFree(Hmm));
+      } else {
+        WARN("Observed error while allocating max (free - 1) memory on ");
+        WARN("GPU: " << i);
+        WARN(" using hipMallocManaged() api with flag 'hipMemAttachGlobal'\n");
+        WARN("Error: " << hipGetErrorString(err));
+        IfTestPassed = false;
+      }
+      err = hipMallocManaged(&Hmm, 1, hipMemAttachHost);
+      if (hipSuccess == err) {
+        HIP_CHECK(hipFree(Hmm));
+      } else {
+        WARN("Observed error while allocating memory size 1 on GPU: " << i);
+        WARN(" with hipMallocManaged() api with flag 'hipMemAttachHost'\n");
+        WARN("Error: " << hipGetErrorString(err));
+        IfTestPassed = false;
+      }
+      err = hipMallocManaged(&Hmm, freeDevMem, hipMemAttachHost);
+      if (hipSuccess == err) {
+        HIP_CHECK(hipFree(Hmm));
+      } else {
+        WARN("Observed error while allocating max free memory on GPU: " << i);
+        WARN(" with hipMallocManaged() api with flag 'hipMemAttachHost'\n");
+        WARN("Error: " << hipGetErrorString(err));
+        IfTestPassed = false;
+      }
+      err = hipMallocManaged(&Hmm, (freeDevMem - 1), hipMemAttachHost);
+      if (hipSuccess == err) {
+        HIP_CHECK(hipFree(Hmm));
+      } else {
+        WARN("Observed error while allocating max (freeDevMem - 1) memory"
+               " on GPU: " << i);
+        WARN(" with hipMallocManaged() api with flag 'hipMemAttachHost'\n");
+        WARN("Error: " << hipGetErrorString(err));
+        IfTestPassed = false;
+      }
+    }
+    REQUIRE(IfTestPassed);
+  } else {
+    SUCCEED("Gpu doesnt support HMM! Hence skipping the test with PASS result");
+  }
+}

tests/catch/unit/deviceLib/CMakeLists.txt

@@ -10,6 +10,7 @@ set(TEST_SRC
     popc.cc
     ldg.cc
     threadfence_system.cc
+    hipTestDeviceSymbol.cc
 )
 
 # skipped for windows compiler issue - Illegal instruction detected

tests/catch/unit/deviceLib/hipTestDeviceSymbol.cc

@@ -0,0 +1,141 @@
+/*
+Copyright (c) 2021 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.
+*/
+
+
+/* Test Case Description: Calling hipMemcpyTo/FromSymbolAsync() using user
+   declared stream obj and hipStreamPerThread*/
+
+#include <hip_test_common.hh>
+#define NUM 1024
+#define SIZE 1024 * 4
+
+__device__ int globalIn[NUM];
+__device__ int globalOut[NUM];
+
+__global__ void Assign(int* Out) {
+    int tid = threadIdx.x + blockIdx.x * blockDim.x;
+    Out[tid] = globalIn[tid];
+    globalOut[tid] = globalIn[tid];
+}
+
+__device__ __constant__ int globalConst[NUM];
+
+__global__ void checkAddress(int* addr, bool* out) {
+    *out = (globalConst == addr);
+}
+
+TEST_CASE("Unit_hipMemcpyToSymbolAsync_ToNFrom") {
+    int *A, *Am, *B, *Ad, *C, *Cm;
+    A = new int[NUM];
+    B = new int[NUM];
+    C = new int[NUM];
+    for (int i = 0; i < NUM; i++) {
+        A[i] = -1 * i;
+        B[i] = 0;
+        C[i] = 0;
+    }
+
+    HIP_CHECK(hipMalloc((void**)&Ad, SIZE));
+    HIP_CHECK(hipHostMalloc((void**)&Am, SIZE));
+    HIP_CHECK(hipHostMalloc((void**)&Cm, SIZE));
+    for (int i = 0; i < NUM; i++) {
+        Am[i] = -1 * i;
+        Cm[i] = 0;
+    }
+
+    hipStream_t stream;
+    HIP_CHECK(hipStreamCreate(&stream));
+    HIP_CHECK(hipMemcpyToSymbolAsync(HIP_SYMBOL(globalIn), Am, SIZE, 0,
+              hipMemcpyHostToDevice, stream));
+    HIP_CHECK(hipStreamSynchronize(stream));
+    hipLaunchKernelGGL(Assign, dim3(1, 1, 1), dim3(NUM, 1, 1), 0, 0, Ad);
+    HIP_CHECK(hipMemcpy(B, Ad, SIZE, hipMemcpyDeviceToHost));
+    HIP_CHECK(hipMemcpyFromSymbolAsync(Cm, HIP_SYMBOL(globalOut), SIZE, 0,
+              hipMemcpyDeviceToHost, stream));
+    HIP_CHECK(hipStreamSynchronize(stream));
+    for (int i = 0; i < NUM; i++) {
+        assert(Am[i] == B[i]);
+        assert(Am[i] == Cm[i]);
+    }
+
+    for (int i = 0; i < NUM; i++) {
+        A[i] = -2 * i;
+        B[i] = 0;
+    }
+
+    HIP_CHECK(hipMemcpyToSymbol(HIP_SYMBOL(globalIn), A, SIZE, 0,
+                                hipMemcpyHostToDevice));
+    hipLaunchKernelGGL(Assign, dim3(1, 1, 1), dim3(NUM, 1, 1), 0, 0, Ad);
+    HIP_CHECK(hipMemcpy(B, Ad, SIZE, hipMemcpyDeviceToHost));
+    HIP_CHECK(hipMemcpyFromSymbol(C, HIP_SYMBOL(globalOut), SIZE, 0,
+                                  hipMemcpyDeviceToHost));
+    for (int i = 0; i < NUM; i++) {
+        assert(A[i] == B[i]);
+        assert(A[i] == C[i]);
+    }
+
+    for (int i = 0; i < NUM; i++) {
+        A[i] = -3 * i;
+        B[i] = 0;
+    }
+    SECTION("Calling hipMemcpyTo/FromSymbol using user declared stream obj") {
+      HIP_CHECK(hipMemcpyToSymbolAsync(HIP_SYMBOL(globalIn), A, SIZE, 0,
+                                        hipMemcpyHostToDevice, stream));
+      HIP_CHECK(hipStreamSynchronize(stream));
+      hipLaunchKernelGGL(Assign, dim3(1, 1, 1), dim3(NUM, 1, 1), 0, 0, Ad);
+      HIP_CHECK(hipMemcpy(B, Ad, SIZE, hipMemcpyDeviceToHost));
+      HIP_CHECK(hipMemcpyFromSymbolAsync(C, HIP_SYMBOL(globalOut), SIZE, 0,
+                                         hipMemcpyDeviceToHost, stream));
+      HIP_CHECK(hipStreamSynchronize(stream));
+    }
+    SECTION("Calling hipMemcpyTo/FromSymbol using hipStreamPerThread") {
+      HIP_CHECK(hipMemcpyToSymbolAsync(HIP_SYMBOL(globalIn), A, SIZE, 0,
+                                        hipMemcpyHostToDevice, hipStreamPerThread));
+      HIP_CHECK(hipStreamSynchronize(hipStreamPerThread));
+      hipLaunchKernelGGL(Assign, dim3(1, 1, 1), dim3(NUM, 1, 1), 0, 0, Ad);
+      HIP_CHECK(hipMemcpy(B, Ad, SIZE, hipMemcpyDeviceToHost));
+      HIP_CHECK(hipMemcpyFromSymbolAsync(C, HIP_SYMBOL(globalOut), SIZE, 0,
+                                         hipMemcpyDeviceToHost, hipStreamPerThread));
+      HIP_CHECK(hipStreamSynchronize(hipStreamPerThread));
+    }
+    for (int i = 0; i < NUM; i++) {
+        assert(A[i] == B[i]);
+        assert(A[i] == C[i]);
+    }
+
+    bool *checkOkD;
+    bool checkOk = false;
+    size_t symbolSize = 0;
+    int *symbolAddress;
+    HIP_CHECK(hipGetSymbolSize(&symbolSize, HIP_SYMBOL(globalConst)));
+    HIP_CHECK(hipGetSymbolAddress((void**) &symbolAddress, HIP_SYMBOL(globalConst)));
+    HIP_CHECK(hipMalloc((void**)&checkOkD, sizeof(bool)));
+    hipLaunchKernelGGL(checkAddress, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0, symbolAddress, checkOkD);
+    HIP_CHECK(hipMemcpy(&checkOk, checkOkD, sizeof(bool), hipMemcpyDeviceToHost));
+    HIP_CHECK(hipFree(checkOkD));
+    HIP_ASSERT(checkOk);
+    HIP_ASSERT((symbolSize == SIZE));
+
+    HIP_CHECK(hipHostFree(Am));
+    HIP_CHECK(hipHostFree(Cm));
+    HIP_CHECK(hipFree(Ad));
+    delete[] A;
+    delete[] B;
+    delete[] C;
+}

tests/catch/unit/graph/CMakeLists.txt

@@ -1,4 +1,4 @@
-# Copyright (c) 2021 Advanced Micro Devices, Inc. All Rights Reserved.
+# Copyright (c) 2022 Advanced Micro Devices, Inc. All Rights Reserved.
 #
 # Permission is hereby granted, free of charge, to any person obtaining a copy
 # of this software and associated documentation files (the "Software"), to deal
@@ -25,6 +25,19 @@ set(TEST_SRC
   hipGraph.cc
   hipSimpleGraphWithKernel.cc
   hipGraphAddMemcpyNode.cc
+  hipGraphClone.cc
+  hipGraphInstantiateWithFlags.cc
+  hipGraphAddHostNode.cc
+  hipGraphAddMemcpyNodeFromSymbol.cc
+  hipGraphChildGraphNodeGetGraph.cc
+  hipGraphNodeFindInClone.cc
+  hipGraphExecHostNodeSetParams.cc
+  hipGraphAddMemcpyNodeToSymbol.cc
+  hipGraphExecMemsetNodeSetParams.cc
+  hipGraphMemcpyNodeSetParamsToSymbol.cc
+  hipGraphDestroyNode.cc
+  hipGraphGetNodes.cc
+  hipGraphGetRootNodes.cc
 )
 
 hip_add_exe_to_target(NAME GraphsTest

tests/catch/unit/graph/hipGraphAddHostNode.cc

@@ -0,0 +1,312 @@
+/*
+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, INCLUDING 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 ANY 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.
+*/
+
+/**
+Testcase Scenarios of hipGraphAddHostNode API:
+
+Functional:
+1. Creates graph, Adds HostNode which updates the variable and validates the result
+2. Create graph, Add Graphnodes and clones the graph. Add Hostnode to the cloned graph
+   and validate the result
+3. Creates graph which performs the square of number in the kernel function and the result
+   is validated in the callback function of hipGraphAddHostNode API
+
+Negative:
+
+1) Pass pGraphNode as nullptr and verify api doesn’t crash, returns error code.
+2) Pass graph as nullptr and verify api doesn’t crash, returns error code.
+3) Pass pNodeParams as nullptr and verify api doesn’t crash, returns error code.
+4) Pass hipHostNodeParams::hipHostFn_t as nullptr and verify api doesn't crash, returns error code.
+*/
+
+#include <hip_test_common.hh>
+#include <hip_test_checkers.hh>
+
+#define SIZE 1024
+
+static int *B_h;
+static int *D_h;
+
+static void callbackfunc(void *A_h) {
+  int *A = reinterpret_cast<int *>(A_h);
+  for (int i = 0; i < SIZE; i++) {
+    A[i] = i;
+  }
+}
+
+static void __global__ vector_square(int *B_d, int *D_d) {
+  for (int i = 0; i < SIZE; i++) {
+    D_d[i] = B_d[i] * B_d[i];
+  }
+}
+static void vectorsquare_callback(void* ptr) {
+  // The callback func is not working with zero parameters
+  // Temporary fix for adding the below 2 lines and ticket
+  // has been raised for the same.
+  int *A = reinterpret_cast<int *>(ptr);
+  A++;
+  for (int i = 0; i < SIZE; i++) {
+    if (D_h[i] != B_h[i] * B_h[i]) {
+      INFO("Validation failed " << D_h[i] << B_h[i]);
+      REQUIRE(false);
+    }
+  }
+}
+/*
+This testcase verifies the negative scenarios of
+hipGraphAddHostNode API
+*/
+TEST_CASE("Unit_hipGraphAddHostNode_Negative") {
+  constexpr size_t N = 1024;
+  hipGraph_t graph;
+  int *A_d{nullptr}, *C_d{nullptr};
+  int *A_h{nullptr}, *C_h{nullptr};
+  HipTest::initArrays<int>(&A_d, nullptr, &C_d,
+                           &A_h, nullptr, &C_h, N, false);
+
+  HIP_CHECK(hipGraphCreate(&graph, 0));
+
+  hipGraphNode_t hostNode;
+  hipHostNodeParams hostParams = {0, 0};
+  hostParams.fn = callbackfunc;
+  hostParams.userData = A_h;
+
+  SECTION("Passing nullptr to graph node") {
+    REQUIRE(hipGraphAddHostNode(nullptr, graph,
+            nullptr,
+            0, &hostParams) == hipErrorInvalidValue);
+  }
+
+  SECTION("Passing nullptr to graph") {
+    REQUIRE(hipGraphAddHostNode(&hostNode, nullptr,
+            nullptr,
+            0, &hostParams) == hipErrorInvalidValue);
+  }
+
+#if HT_NVIDIA
+  SECTION("Passing nullptr to host params") {
+    REQUIRE(hipGraphAddHostNode(&hostNode, graph,
+            nullptr,
+            0, nullptr) == hipErrorInvalidValue);
+  }
+#endif
+
+  SECTION("Passing nullptr to host func") {
+    hostParams.fn = nullptr;
+    REQUIRE(hipGraphAddHostNode(&hostNode, graph,
+            nullptr,
+            0, &hostParams) == hipErrorInvalidValue);
+  }
+
+  HIP_CHECK(hipGraphDestroy(graph));
+}
+/*
+This testcase verifies hipGraphAddHostNode API in cloned graph
+Creates graph, Add graph nodes and clone the graph
+Add HostNode to the cloned graph and validate the result
+*/
+TEST_CASE("Unit_hipGraphAddHostNode_ClonedGraphwithHostNode") {
+  constexpr size_t N = 1024;
+  constexpr size_t Nbytes = N * sizeof(int);
+  hipGraph_t graph;
+  hipGraphExec_t graphExec;
+  int *A_d{nullptr}, *C_d{nullptr};
+  int *A_h{nullptr}, *C_h{nullptr};
+  HipTest::initArrays<int>(&A_d, nullptr, &C_d,
+                           &A_h, nullptr, &C_h, N, false);
+
+  HIP_CHECK(hipGraphCreate(&graph, 0));
+  hipGraphNode_t memcpyH2D_A, memcpyH2D_C,
+                 memcpyD2H_AC;
+  hipStream_t streamForGraph;
+  HIP_CHECK(hipStreamCreate(&streamForGraph));
+  HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr,
+                                    0, A_d, A_h,
+                                    Nbytes, hipMemcpyHostToDevice));
+  HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_C, graph, nullptr,
+                                    0, C_d, C_h,
+                                    Nbytes, hipMemcpyHostToDevice));
+  HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_AC, graph, nullptr,
+                                    0, A_h, C_d,
+                                    Nbytes, hipMemcpyDeviceToHost));
+
+  hipGraph_t clonedgraph;
+  HIP_CHECK(hipGraphClone(&clonedgraph, graph));
+
+  hipGraphNode_t hostNode;
+  hipHostNodeParams hostParams = {0, 0};
+  hostParams.fn = callbackfunc;
+  hostParams.userData = A_h;
+  HIP_CHECK(hipGraphAddHostNode(&hostNode, clonedgraph,
+                                 nullptr,
+                                 0, &hostParams));
+
+  HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_A,
+                                    &memcpyD2H_AC, 1));
+  HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_C,
+                                    &memcpyD2H_AC, 1));
+
+  // Instantiate and launch the cloned graph
+  HIP_CHECK(hipGraphInstantiate(&graphExec, clonedgraph, nullptr, nullptr, 0));
+  HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
+  HIP_CHECK(hipStreamSynchronize(streamForGraph));
+
+  // Verify execution result
+  for (size_t i = 0; i < N; i++) {
+    if (A_h[i] != static_cast<int>(i)) {
+      INFO("Validation failed i " <<  i  << "C_h[i] "<< C_h[i]);
+      REQUIRE(false);
+    }
+  }
+
+  HipTest::freeArrays<int>(A_d, nullptr, C_d, A_h, nullptr, C_h, false);
+  HIP_CHECK(hipGraphExecDestroy(graphExec));
+  HIP_CHECK(hipGraphDestroy(graph));
+  HIP_CHECK(hipGraphDestroy(clonedgraph));
+  HIP_CHECK(hipStreamDestroy(streamForGraph));
+}
+
+/*
+This testcase verifies the square of number by
+creating graph, Add kernel node which does the square
+of number and the result is validated byhipGrahAddHostNode API
+*/
+TEST_CASE("Unit_hipGraphAddHostNode_VectorSquare") {
+  constexpr size_t N = 1024;
+  constexpr size_t Nbytes = N * sizeof(int);
+  hipGraph_t graph;
+  hipGraphExec_t graphExec;
+  int *A_d{nullptr}, *A_h{nullptr}, *B_d{nullptr}, *D_d{nullptr};
+  int *param = reinterpret_cast<int *>(sizeof(int));;
+  HipTest::initArrays<int>(&A_d, &B_d, &D_d,
+                           &A_h, &B_h, &D_h, N, false);
+  HIP_CHECK(hipGraphCreate(&graph, 0));
+  hipGraphNode_t memcpyH2D_B, memcpyH2D_D, memcpyD2H_D, kernel_vecAdd;
+  hipKernelNodeParams kernelNodeParams{};
+  hipStream_t streamForGraph;
+  HIP_CHECK(hipStreamCreate(&streamForGraph));
+  hipGraphNode_t hostNode;
+  hipHostNodeParams hostParams = {0, 0};
+  hostParams.fn = vectorsquare_callback;
+  hostParams.userData = param;
+
+  HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_B, graph, nullptr,
+                                    0, B_d, B_h,
+                                    Nbytes, hipMemcpyHostToDevice));
+  HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_D, graph, nullptr,
+                                    0, D_d, D_h,
+                                    Nbytes, hipMemcpyHostToDevice));
+
+  void* kernelArgs2[] = {&B_d, &D_d};
+  kernelNodeParams.func = reinterpret_cast<void *>(vector_square);
+  kernelNodeParams.gridDim = dim3(1);
+  kernelNodeParams.blockDim = dim3(1);
+  kernelNodeParams.sharedMemBytes = 0;
+  kernelNodeParams.kernelParams = reinterpret_cast<void**>(kernelArgs2);
+  kernelNodeParams.extra = nullptr;
+  HIP_CHECK(hipGraphAddKernelNode(&kernel_vecAdd, graph, nullptr, 0,
+                                  &kernelNodeParams));
+  HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_D, graph, nullptr,
+                                    0, D_h, D_d,
+                                    Nbytes, hipMemcpyDeviceToHost));
+
+  HIP_CHECK(hipGraphAddHostNode(&hostNode, graph,
+                                 nullptr,
+                                 0, &hostParams));
+  HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_B, &kernel_vecAdd,
+                                    1));
+  HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_D, &kernel_vecAdd,
+                                    1));
+  HIP_CHECK(hipGraphAddDependencies(graph, &kernel_vecAdd,
+                                    &memcpyD2H_D, 1));
+  HIP_CHECK(hipGraphAddDependencies(graph, &memcpyD2H_D,
+                                    &hostNode, 1));
+
+  // Instantiate and launch the graph
+  HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
+  HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
+  HIP_CHECK(hipStreamSynchronize(streamForGraph));
+
+  HipTest::freeArrays<int>(A_d, B_d, D_d, A_h, B_h, D_h, false);
+  HIP_CHECK(hipGraphExecDestroy(graphExec));
+  HIP_CHECK(hipGraphDestroy(graph));
+  HIP_CHECK(hipStreamDestroy(streamForGraph));
+}
+/*
+This testcase verifies the following scenario
+Create graph, calls the host function and updates
+the parameters in the callback function and
+validates it.
+*/
+TEST_CASE("Unit_hipGraphAddHostNode_BasicFunc") {
+  constexpr size_t N = 1024;
+  constexpr size_t Nbytes = N * sizeof(int);
+  hipGraph_t graph;
+  hipGraphExec_t graphExec;
+  int *A_d{nullptr}, *C_d{nullptr};
+  int *A_h{nullptr}, *C_h{nullptr};
+  HipTest::initArrays<int>(&A_d, nullptr, &C_d,
+                           &A_h, nullptr, &C_h, N, false);
+
+  HIP_CHECK(hipGraphCreate(&graph, 0));
+  hipGraphNode_t memcpyH2D_A, memcpyD2H_AC, memcpyH2D_C;
+  hipStream_t streamForGraph;
+  HIP_CHECK(hipStreamCreate(&streamForGraph));
+  HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr,
+                                    0, A_d, A_h,
+                                    Nbytes, hipMemcpyHostToDevice));
+  HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_C, graph, nullptr,
+                                    0, C_d, C_h,
+                                    Nbytes, hipMemcpyHostToDevice));
+  HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_AC, graph, nullptr,
+                                    0, A_h, C_d,
+                                    Nbytes, hipMemcpyDeviceToHost));
+  hipGraphNode_t hostNode;
+  hipHostNodeParams hostParams = {0, 0};
+  hostParams.fn = callbackfunc;
+  hostParams.userData = A_h;
+  HIP_CHECK(hipGraphAddHostNode(&hostNode, graph,
+                                 nullptr,
+                                 0, &hostParams));
+
+  HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_A,
+                                    &memcpyD2H_AC, 1));
+  HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_C,
+                                    &memcpyD2H_AC, 1));
+  HIP_CHECK(hipGraphAddDependencies(graph, &memcpyD2H_AC,
+                                    &hostNode, 1));
+
+  // Instantiate and launch the graph
+  HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
+  HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
+  HIP_CHECK(hipStreamSynchronize(streamForGraph));
+
+  // Verify execution result
+  for (size_t i = 0; i < N; i++) {
+    if (A_h[i] != static_cast<int>(i)) {
+      INFO("Validation failed i " <<  i  << "A_h[i] "<< A_h[i]);
+      REQUIRE(false);
+    }
+  }
+
+  HipTest::freeArrays<int>(A_d, nullptr, C_d, A_h, nullptr, C_h, false);
+  HIP_CHECK(hipGraphExecDestroy(graphExec));
+  HIP_CHECK(hipGraphDestroy(graph));
+  HIP_CHECK(hipStreamDestroy(streamForGraph));
+}

tests/catch/unit/graph/hipGraphAddMemcpyNode.cc

@@ -19,18 +19,22 @@ THE SOFTWARE.
 
 /**
 Testcase Scenarios :
- 1) Add multiple Memcpy nodes to graph and verify node execution is
- working as expected.
+1) Add memcpy node to graph and verify memcpy operation is success for all memcpy kinds(H2D, D2H and D2D).
+Memcpy nodes are added and assigned to default device.
+2) Perform memcpy operation for 1D, 2D and 3D arrays on default device and verify the results.
+3) Add memcpy node to graph and verify memcpy operation is success for all memcpy kinds(H2D, D2H and D2D).
+Memcpy nodes are added and assigned to Peer device.
+4) Perform memcpy operation for 1D, 2D and 3D arrays on Peer device and verify the results.
 */
 
 #include <hip_test_common.hh>
 #include <hip_test_checkers.hh>
 
-/**
- * Functional Test adds memcpy nodes of types H2D, D2D and D2H to graph
- * and verifies execution sequence by launching graph.
- */
-TEST_CASE("Unit_hipGraphAddMemcpyNode_Functional") {
+#define ZSIZE 32
+#define YSIZE 32
+#define XSIZE 32
+
+void validateMemcpyNode3DArray(bool peerAccess = false) {
   constexpr int width{10}, height{10}, depth{10};
   hipArray *devArray1, *devArray2;
   hipChannelFormatKind formatKind = hipChannelFormatKindSigned;
@@ -42,6 +46,7 @@ TEST_CASE("Unit_hipGraphAddMemcpyNode_Functional") {
   hipStream_t streamForGraph;
   hipGraphExec_t graphExec;
 
+  HIP_CHECK(hipSetDevice(0));
   int *hData = reinterpret_cast<int*>(malloc(size));
   int *hOutputData = reinterpret_cast<int *>(malloc(size));
 
@@ -69,6 +74,12 @@ TEST_CASE("Unit_hipGraphAddMemcpyNode_Functional") {
                        make_hipExtent(width, height, depth), hipArrayDefault));
   HIP_CHECK(hipGraphCreate(&graph, 0));
 
+  // For peer access test, Memory is allocated on device(0)
+  // while memcpy nodes are allocated and assigned to peer device(1)
+  if (peerAccess) {
+    HIP_CHECK(hipSetDevice(1));
+  }
+
   // Host to Device
   memset(&myparams, 0x0, sizeof(hipMemcpy3DParms));
   myparams.srcPos = make_hipPos(0, 0, 0);
@@ -79,7 +90,6 @@ TEST_CASE("Unit_hipGraphAddMemcpyNode_Functional") {
   myparams.dstArray = devArray1;
   myparams.kind = hipMemcpyHostToDevice;
 
-
   HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, nullptr, 0, &myparams));
   dependencies.push_back(memcpyNode);
 
@@ -126,3 +136,247 @@ TEST_CASE("Unit_hipGraphAddMemcpyNode_Functional") {
   free(hData);
   free(hOutputData);
 }
+
+void validateMemcpyNode2DArray(bool peerAccess = false) {
+  int harray2D[YSIZE][XSIZE]{};
+  int harray2Dres[YSIZE][XSIZE]{};
+  constexpr int width{XSIZE}, height{YSIZE};
+  hipArray *devArray1, *devArray2;
+  hipChannelFormatKind formatKind = hipChannelFormatKindSigned;
+  hipMemcpy3DParms myparams;
+  hipGraph_t graph;
+  hipGraphNode_t memcpyNode;
+  std::vector<hipGraphNode_t> dependencies;
+  hipStream_t streamForGraph;
+  hipGraphExec_t graphExec;
+
+  HIP_CHECK(hipSetDevice(0));
+  HIP_CHECK(hipStreamCreate(&streamForGraph));
+  // Initialize 2D object
+  for (int i = 0; i < YSIZE; i++) {
+    for (int j = 0; j < XSIZE; j++) {
+      harray2D[i][j] = i + j + 1;
+    }
+  }
+
+  hipChannelFormatDesc channelDesc = hipCreateChannelDesc(sizeof(int)*8,
+                                                          0, 0, 0, formatKind);
+  // Allocate 2D device array by passing depth(0)
+  HIP_CHECK(hipMalloc3DArray(&devArray1, &channelDesc,
+                       make_hipExtent(width, height, 0), hipArrayDefault));
+  HIP_CHECK(hipMalloc3DArray(&devArray2, &channelDesc,
+                       make_hipExtent(width, height, 0), hipArrayDefault));
+  HIP_CHECK(hipGraphCreate(&graph, 0));
+
+  // For peer access test, Memory is allocated on device(0)
+  // while memcpy nodes are allocated and assigned to peer device(1)
+  if (peerAccess) {
+    HIP_CHECK(hipSetDevice(1));
+  }
+
+  // Host to Device
+  memset(&myparams, 0x0, sizeof(hipMemcpy3DParms));
+  myparams.srcPos = make_hipPos(0, 0, 0);
+  myparams.dstPos = make_hipPos(0, 0, 0);
+  myparams.extent = make_hipExtent(width, height, 1);
+  myparams.srcPtr = make_hipPitchedPtr(harray2D, width * sizeof(int),
+                                      width, height);
+  myparams.dstArray = devArray1;
+  myparams.kind = hipMemcpyHostToDevice;
+
+  HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, nullptr, 0, &myparams));
+  dependencies.push_back(memcpyNode);
+
+  // Device to Device
+  memset(&myparams, 0x0, sizeof(hipMemcpy3DParms));
+  myparams.srcPos = make_hipPos(0, 0, 0);
+  myparams.dstPos = make_hipPos(0, 0, 0);
+  myparams.srcArray = devArray1;
+  myparams.dstArray = devArray2;
+  myparams.extent = make_hipExtent(width, height, 1);
+  myparams.kind = hipMemcpyDeviceToDevice;
+
+  HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, dependencies.data(),
+                                             dependencies.size(), &myparams));
+  dependencies.clear();
+  dependencies.push_back(memcpyNode);
+
+  // Device to host
+  memset(&myparams, 0x0, sizeof(hipMemcpy3DParms));
+  myparams.srcPos = make_hipPos(0, 0, 0);
+  myparams.dstPos = make_hipPos(0, 0, 0);
+  myparams.extent = make_hipExtent(width, height, 1);
+  myparams.dstPtr = make_hipPitchedPtr(harray2Dres, width * sizeof(int),
+                                      width, height);
+  myparams.srcArray = devArray2;
+  myparams.kind = hipMemcpyDeviceToHost;
+
+  HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, dependencies.data(),
+                                             dependencies.size(), &myparams));
+
+  // Instantiate and launch the graph
+  HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
+  HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
+  HIP_CHECK(hipStreamSynchronize(streamForGraph));
+
+  // Validate result
+  for (int i = 0; i < YSIZE; i++) {
+    for (int j = 0; j < XSIZE; j++) {
+      if (harray2D[i][j] != harray2Dres[i][j]) {
+        INFO("harray2D: " << harray2D[i][j] << "harray2Dres: "
+              << harray2Dres[i][j] << " mismatch at (i,j) : " << i << j);
+        REQUIRE(false);
+      }
+    }
+  }
+  HIP_CHECK(hipGraphExecDestroy(graphExec));
+  HIP_CHECK(hipGraphDestroy(graph));
+  HIP_CHECK(hipStreamDestroy(streamForGraph));
+  hipFreeArray(devArray1);
+  hipFreeArray(devArray2);
+}
+
+void validateMemcpyNode1DArray(bool peerAccess = false) {
+  int harray1D[XSIZE]{};
+  int harray1Dres[XSIZE]{};
+  constexpr int width{XSIZE};
+  hipArray *devArray1, *devArray2;
+  hipChannelFormatKind formatKind = hipChannelFormatKindSigned;
+  hipMemcpy3DParms myparams;
+  hipGraph_t graph;
+  hipGraphNode_t memcpyNode;
+  std::vector<hipGraphNode_t> dependencies;
+  hipStream_t streamForGraph;
+  hipGraphExec_t graphExec;
+
+  HIP_CHECK(hipSetDevice(0));
+  HIP_CHECK(hipStreamCreate(&streamForGraph));
+  // Initialize 1D object
+  for (int i = 0; i < XSIZE; i++) {
+    harray1D[i] = i + 1;
+  }
+
+  hipChannelFormatDesc channelDesc = hipCreateChannelDesc(sizeof(int)*8,
+                                                          0, 0, 0, formatKind);
+  // Allocate 1D device array by passing depth(0), height(0)
+  HIP_CHECK(hipMalloc3DArray(&devArray1, &channelDesc,
+                       make_hipExtent(width, 0, 0), hipArrayDefault));
+  HIP_CHECK(hipMalloc3DArray(&devArray2, &channelDesc,
+                       make_hipExtent(width, 0, 0), hipArrayDefault));
+  HIP_CHECK(hipGraphCreate(&graph, 0));
+
+  // For peer access test, Memory is allocated on device(0)
+  // while memcpy nodes are allocated and assigned to peer device(1)
+  if (peerAccess) {
+    HIP_CHECK(hipSetDevice(1));
+  }
+
+  // Host to Device
+  memset(&myparams, 0x0, sizeof(hipMemcpy3DParms));
+  myparams.srcPos = make_hipPos(0, 0, 0);
+  myparams.dstPos = make_hipPos(0, 0, 0);
+  myparams.extent = make_hipExtent(width, 1, 1);
+  myparams.srcPtr = make_hipPitchedPtr(harray1D, width * sizeof(int),
+                                      width, 1);
+  myparams.dstArray = devArray1;
+  myparams.kind = hipMemcpyHostToDevice;
+
+  HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, nullptr, 0, &myparams));
+  dependencies.push_back(memcpyNode);
+
+  // Device to Device
+  memset(&myparams, 0x0, sizeof(hipMemcpy3DParms));
+  myparams.srcPos = make_hipPos(0, 0, 0);
+  myparams.dstPos = make_hipPos(0, 0, 0);
+  myparams.srcArray = devArray1;
+  myparams.dstArray = devArray2;
+  myparams.extent = make_hipExtent(width, 1, 1);
+  myparams.kind = hipMemcpyDeviceToDevice;
+
+  HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, dependencies.data(),
+                                             dependencies.size(), &myparams));
+  dependencies.clear();
+  dependencies.push_back(memcpyNode);
+
+  // Device to host
+  memset(&myparams, 0x0, sizeof(hipMemcpy3DParms));
+  myparams.srcPos = make_hipPos(0, 0, 0);
+  myparams.dstPos = make_hipPos(0, 0, 0);
+  myparams.extent = make_hipExtent(width, 1, 1);
+  myparams.dstPtr = make_hipPitchedPtr(harray1Dres, width * sizeof(int),
+                                      width, 1);
+  myparams.srcArray = devArray2;
+  myparams.kind = hipMemcpyDeviceToHost;
+
+  HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, dependencies.data(),
+                                              dependencies.size(), &myparams));
+
+  // Instantiate and launch the graph
+  HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
+  HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
+  HIP_CHECK(hipStreamSynchronize(streamForGraph));
+
+  // Validate result
+  for (int i = 0; i < XSIZE; i++) {
+    if (harray1D[i] != harray1Dres[i]) {
+      INFO("harray1D: " << harray1D[i] << " harray1Dres: " << harray1Dres[i]
+            << " mismatch at : " << i);
+      REQUIRE(false);
+    }
+  }
+  HIP_CHECK(hipGraphExecDestroy(graphExec));
+  HIP_CHECK(hipGraphDestroy(graph));
+  HIP_CHECK(hipStreamDestroy(streamForGraph));
+  hipFreeArray(devArray1);
+  hipFreeArray(devArray2);
+}
+
+/**
+ * Basic Functional Tests adds memcpy nodes of types H2D, D2D and D2H to graph
+ * and verifies execution sequence by launching graph on default device.
+ * Tests also verify memcpy node addition with 1D, 2D and 3D objects.
+ */
+TEST_CASE("Unit_hipGraphAddMemcpyNode_BasicFunctional") {
+  SECTION("Memcpy with 3D array on default device") {
+    validateMemcpyNode3DArray();
+  }
+
+  SECTION("Memcpy with 2D array on default device") {
+    validateMemcpyNode2DArray();
+  }
+
+  SECTION("Memcpy with 1D array on default device") {
+    validateMemcpyNode1DArray();
+  }
+}
+
+/**
+ * Peer access tests adds and assigns memcpy nodes of types H2D, D2D and D2H
+ * to peer device. Memory allocations happen on device(0) and memcpy operations
+ * are performed from device(1).
+ * Tests also verify memcpy node addition with 1D, 2D and 3D objects.
+ */
+TEST_CASE("Unit_hipGraphAddMemcpyNode_PeerAccessFunctional") {
+  int numDevices{}, peerAccess{};
+  HIP_CHECK(hipGetDeviceCount(&numDevices));
+  if (numDevices > 1) {
+    HIP_CHECK(hipDeviceCanAccessPeer(&peerAccess, 1, 0));
+  }
+
+  if (!peerAccess) {
+    WARN("Skipping test as peer device access is not found!");
+    return;
+  }
+
+  SECTION("Memcpy with 3D array on peer device") {
+    validateMemcpyNode3DArray(true);
+  }
+
+  SECTION("Memcpy with 2D array on peer device") {
+    validateMemcpyNode2DArray(true);
+  }
+
+  SECTION("Memcpy with 1D array on peer device") {
+    validateMemcpyNode1DArray(true);
+  }
+}

tests/catch/unit/graph/hipGraphAddMemcpyNodeFromSymbol.cc

@@ -0,0 +1,435 @@
+/*
+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, INCLUDING 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 ANY 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.
+*/
+
+/**
+Testcase Scenarios of hipGraphAddMemcpyNodeFromSymbol API:
+
+Functional :
+
+1. Allocate global symbol memory, add the MemcpyNodeFromSymbol
+   node to the graph and verify  for different memory kinds
+2. Allocate const memory  add the MemcpyNodeFromSymbol node to
+   the graph and verify  for different memory kinds
+3. Allocate global symbol memory and  device memory in GPU-0
+   and perform MemcpyToSymbol from peer GPU by adding it to the graph node.
+4. Allocate const symbol memory and  device memory in GPU-0
+   and perform MemcpyToSymbol from peer GPU by adding it to the graph node.
+5. Allocate global memory, Add MemcpyFromSymbolNode,KernelNode and memcpynode and validating
+   the behaviour
+
+Negative :
+
+1) Pass nullptr to graph node
+2) Pass nullptr to graph
+3) Pass nullptr to dependencies
+4) Pass invalid numDependencies
+5) Pass nullptr to dst
+6) Pass nullptr to symbol
+7) Pass invalid count 
+8) Pass offset+count greater than allocated size
+9) Pass unintialized graph
+*/
+
+#include <hip_test_common.hh>
+#include <hip_test_checkers.hh>
+#include <limits>
+#define SIZE 256
+
+__device__ int globalIn[SIZE];
+__device__ int globalOut[SIZE];
+__device__ __constant__ int globalConst[SIZE];
+
+__global__ void MemcpyFromSymbolKernel(int* B_d) {
+  for (int i = 0 ; i < SIZE; i++) {
+    globalIn[i] = B_d[i];
+  }
+}
+
+/* This testcase verifies negative scenarios of
+   hipGraphAddMemcpyNodeFromSymbol API */
+TEST_CASE("Unit_hipGraphAddMemcpyNodeFromSymbol_Negative") {
+  constexpr size_t Nbytes = SIZE * sizeof(int);
+  int *A_d{nullptr}, *B_d{nullptr};
+  int *A_h{nullptr}, *B_h{nullptr};
+  HipTest::initArrays<int>(&A_d, &B_d, nullptr,
+                           &A_h, &B_h, nullptr, SIZE, false);
+
+  hipGraph_t graph;
+  hipGraphNode_t memcpyToSymbolNode, memcpyH2D_A;
+  std::vector<hipGraphNode_t> dependencies;
+  HIP_CHECK(hipGraphCreate(&graph, 0));
+
+  // Adding MemcpyNode
+  HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, 0, A_d, A_h,
+                                    Nbytes, hipMemcpyHostToDevice));
+  dependencies.push_back(memcpyH2D_A);
+
+  // Adding MemcpyNodeToSymbol
+  HIP_CHECK(hipGraphAddMemcpyNodeToSymbol(&memcpyToSymbolNode, graph,
+                                            dependencies.data(),
+                                            dependencies.size(),
+                                            HIP_SYMBOL(globalIn),
+                                            A_d, Nbytes, 0,
+                                            hipMemcpyDeviceToDevice));
+  dependencies.clear();
+  dependencies.push_back(memcpyToSymbolNode);
+
+#if HT_NVIDIA
+  hipGraphNode_t memcpyFromSymbolNode;
+  SECTION("Passing nullptr to graph") {
+    REQUIRE(hipGraphAddMemcpyNodeFromSymbol(&memcpyFromSymbolNode, nullptr,
+                                            dependencies.data(),
+                                            dependencies.size(),
+                                            B_d,
+                                            HIP_SYMBOL(globalIn),
+                                            Nbytes, 0,
+                                            hipMemcpyDeviceToDevice)
+                                            == hipErrorInvalidValue);
+  }
+
+  SECTION("Passing nullptr to graph node") {
+    REQUIRE(hipGraphAddMemcpyNodeFromSymbol(nullptr, graph,
+                                          dependencies.data(),
+                                          dependencies.size(),
+                                          B_d,
+                                          HIP_SYMBOL(globalIn),
+                                          Nbytes, 0,
+                                          hipMemcpyDeviceToDevice)
+                                          == hipErrorInvalidValue);
+  }
+
+  SECTION("Passing size > 1 and dependencies as nullptr") {
+    REQUIRE(hipGraphAddMemcpyNodeFromSymbol(&memcpyFromSymbolNode, graph,
+                                          nullptr,
+                                          1,
+                                          B_d,
+                                          HIP_SYMBOL(globalIn),
+                                          Nbytes, 0,
+                                          hipMemcpyDeviceToDevice)
+                                          == hipErrorInvalidValue);
+  }
+
+  SECTION("Passing invalid dependencies size") {
+    REQUIRE(hipGraphAddMemcpyNodeFromSymbol(&memcpyFromSymbolNode, graph,
+                                          dependencies.data(),
+                                          10,
+                                          B_d,
+                                          HIP_SYMBOL(globalIn),
+                                          Nbytes, 0,
+                                          hipMemcpyDeviceToDevice)
+                                          == hipErrorInvalidValue);
+  }
+
+  SECTION("Passing nullptr to dst") {
+    REQUIRE(hipGraphAddMemcpyNodeFromSymbol(&memcpyFromSymbolNode, graph,
+                                          dependencies.data(),
+                                          dependencies.size(),
+                                          nullptr,
+                                          HIP_SYMBOL(globalIn), Nbytes, 0,
+                                          hipMemcpyDeviceToDevice)
+                                          == hipErrorInvalidValue);
+  }
+
+  SECTION("Passing nullptr to source") {
+    REQUIRE(hipGraphAddMemcpyNodeFromSymbol(&memcpyFromSymbolNode, graph,
+                                          dependencies.data(),
+                                          dependencies.size(),
+                                          B_d,
+                                          nullptr, Nbytes, 0,
+                                          hipMemcpyDeviceToDevice)
+                                          == hipErrorInvalidSymbol);
+  }
+
+  SECTION("Passing offset+size > max size") {
+    REQUIRE(hipGraphAddMemcpyNodeFromSymbol(&memcpyFromSymbolNode, graph,
+                                          dependencies.data(),
+                                          dependencies.size(),
+                                          B_d,
+                                          HIP_SYMBOL(globalIn),
+                                          Nbytes, 10,
+                                          hipMemcpyDeviceToDevice)
+                                          == hipErrorInvalidValue);
+  }
+
+  SECTION("Passing Max count") {
+  REQUIRE(hipGraphAddMemcpyNodeFromSymbol(&memcpyFromSymbolNode, graph,
+                                        dependencies.data(),
+                                        dependencies.size(),
+                                        B_d,
+                                        HIP_SYMBOL(globalIn),
+                                        std::numeric_limits<int>::max(), 0,
+                                        hipMemcpyDeviceToDevice)
+                                        == hipErrorInvalidValue);
+  }
+
+  SECTION("Pass Unintialized graph") {
+  hipGraph_t unint_graph;
+  REQUIRE(hipGraphAddMemcpyNodeFromSymbol(&memcpyFromSymbolNode, unint_graph,
+                                        dependencies.data(),
+                                        dependencies.size(),
+                                        B_d,
+                                        HIP_SYMBOL(globalIn),
+                                        Nbytes, 0,
+                                        hipMemcpyDeviceToDevice)
+                                        == hipErrorInvalidValue);
+  }
+#endif
+
+  HipTest::freeArrays<int>(A_d, B_d, nullptr,
+                      A_h, B_h, nullptr, false);
+  HIP_CHECK(hipGraphDestroy(graph));
+}
+/*
+This function is used to verify the following scenarios
+1. Create global variable, allocate Memory in GPU-0 and create dependency graph of
+   hipGraphAddMemcpyNodeFromSymbol API in GPU-1 and validate the result
+2. Allocate global memory, Create dependency graph and validate the result on GPU-0
+3. Allocate global const memory, Create dependency graph and validate the result on GPU-0
+4. Create global const variable, allocate Memory in GPU-0 and create dependency graph of
+   hipGraphAddMemcpyNodeFromSymbol API in GPU-1 and validate the result
+*/
+
+void hipGraphAddMemcpyNodeFromSymbol_GlobalMemory(bool device_ctxchg = false,
+                                                  bool const_device_var =
+                                                  false) {
+  constexpr size_t Nbytes = SIZE * sizeof(int);
+  int *A_d{nullptr};
+  int *A_h{nullptr}, *B_h{nullptr};
+  HipTest::initArrays<int>(&A_d, nullptr, nullptr,
+                           &A_h, &B_h, nullptr, SIZE, false);
+
+  hipGraph_t graph;
+  hipGraphExec_t graphExec;
+  hipGraphNode_t memcpyToSymbolNode, memcpyFromSymbolNode, memcpyH2D_A;
+  std::vector<hipGraphNode_t> dependencies;
+  HIP_CHECK(hipGraphCreate(&graph, 0));
+
+  if (device_ctxchg) {
+    HIP_CHECK(hipSetDevice(1));
+    HIP_CHECK(hipDeviceEnablePeerAccess(0, 0));
+  }
+  // Adding MemcpyNode
+  HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, 0, A_d, A_h,
+                                    Nbytes, hipMemcpyHostToDevice));
+  dependencies.push_back(memcpyH2D_A);
+
+  // Adding MemcpyNodeToSymbol
+  if (const_device_var) {
+    HIP_CHECK(hipGraphAddMemcpyNodeToSymbol(&memcpyToSymbolNode, graph,
+          dependencies.data(),
+          dependencies.size(),
+          HIP_SYMBOL(globalConst),
+          A_d, Nbytes, 0,
+          hipMemcpyDeviceToDevice));
+
+  } else {
+    HIP_CHECK(hipGraphAddMemcpyNodeToSymbol(&memcpyToSymbolNode, graph,
+          dependencies.data(),
+          dependencies.size(),
+          HIP_SYMBOL(globalIn),
+          A_d, Nbytes, 0,
+          hipMemcpyDeviceToDevice));
+  }
+  dependencies.clear();
+  dependencies.push_back(memcpyToSymbolNode);
+
+
+  // Adding MemcpyNodeFromSymbol
+  if (const_device_var) {
+    HIP_CHECK(hipGraphAddMemcpyNodeFromSymbol(&memcpyFromSymbolNode, graph,
+          dependencies.data(),
+          dependencies.size(),
+          B_h,
+          HIP_SYMBOL(globalConst),
+          Nbytes, 0,
+          hipMemcpyDeviceToHost));
+  } else {
+    HIP_CHECK(hipGraphAddMemcpyNodeFromSymbol(&memcpyFromSymbolNode, graph,
+          dependencies.data(),
+          dependencies.size(),
+          B_h,
+          HIP_SYMBOL(globalIn),
+          Nbytes, 0,
+          hipMemcpyDeviceToHost));
+  }
+
+
+  // Instantiate and launch the graph
+  HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
+  HIP_CHECK(hipGraphLaunch(graphExec, 0));
+
+  // Validating the result
+  for (int i = 0; i < SIZE; i++) {
+    if (B_h[i] != A_h[i]) {
+       WARN("Validation failed B_h[i] " << B_h[i] << "A_h[i] " << A_h[i]);
+       REQUIRE(false);
+    }
+  }
+
+  HipTest::freeArrays<int>(A_d, nullptr, nullptr,
+                      A_h, B_h, nullptr, false);
+  HIP_CHECK(hipGraphExecDestroy(graphExec));
+  HIP_CHECK(hipGraphDestroy(graph));
+}
+/*
+This testcase verifies allocating global symbol memory,
+add the MemcpyNodeFromSymbol node to the graph and
+erifying the result
+*/
+TEST_CASE("Unit_hipGraphAddMemcpyNodeFromSymbol_GlobalMemory") {
+  hipGraphAddMemcpyNodeFromSymbol_GlobalMemory(false, false);
+}
+
+/*
+This testcase verifies allocating global const symbol memory,
+add the MemcpyNodeFromSymbol node to the graph and
+verifying the result
+*/
+
+TEST_CASE("Unit_hipGraphAddMemcpyNodeFromSymbol_GlobalConstMemory") {
+  hipGraphAddMemcpyNodeFromSymbol_GlobalMemory(false, true);
+}
+
+/*
+This testcase verifies allocating global symbol memory and device variables
+in GPU-0 and add the MemcpyNodeFromSymbol node to the graph and
+verifying the result in GPU-1
+*/
+#if HT_NVIDIA
+TEST_CASE("Unit_hipGraphAddMemcpyNodeFromSymbol_GlobalMemoryPeerDevice") {
+  int numDevices = 0;
+  int canAccessPeer = 0;
+  if (numDevices > 1) {
+    HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1));
+    if (canAccessPeer) {
+      hipGraphAddMemcpyNodeFromSymbol_GlobalMemory(true, false);
+    } else {
+      SUCCEED("Machine does not seem to have P2P");
+    }
+  } else {
+    SUCCEED("skipped the testcase as no of devices is less than 2");
+  }
+}
+
+/*
+This testcase verifies allocating global const symbol memory and device variables
+in GPU-0 and add the MemcpyNodeFromSymbol node to the graph and
+verifying the result in GPU-1
+*/
+
+TEST_CASE("Unit_hipGraphAddMemcpyNodeFromSymbol_GlobalConstMemoryPeerDevice") {
+  int numDevices = 0;
+  int canAccessPeer = 0;
+  if (numDevices > 1) {
+    HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1));
+    if (canAccessPeer) {
+      hipGraphAddMemcpyNodeFromSymbol_GlobalMemory(true, true);
+    } else {
+      SUCCEED("Machine does not seem to have P2P");
+    }
+  } else {
+    SUCCEED("skipped the testcase as no of devices is less than 2");
+  }
+}
+#endif
+/*
+This testcaser verifies allocating global memory,
+Add MemcpyFromSymbolNode,KernelNode and memcpynode and validating
+the behaviour
+*/
+TEST_CASE("Unit_hipGraphAddMemcpyNodeFromSymbol_GlobalMemoryWithKernel") {
+  constexpr size_t Nbytes = SIZE * sizeof(int);
+  constexpr auto blocksPerCU = 6;  // to hide latency
+  constexpr auto threadsPerBlock = 256;
+  unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, SIZE);
+  hipGraphNode_t memcpyfromsymbolkernel, memcpyD2H_B;
+  hipKernelNodeParams kernelNodeParams{};
+  int *A_d{nullptr}, *B_d{nullptr};
+  int *A_h{nullptr}, *B_h{nullptr};
+  HipTest::initArrays<int>(&A_d, &B_d, nullptr,
+                           &A_h, &B_h, nullptr, SIZE, false);
+
+  hipGraph_t graph;
+  hipGraphExec_t graphExec;
+  hipGraphNode_t memcpyToSymbolNode, memcpyFromSymbolNode, memcpyH2D_A;
+  std::vector<hipGraphNode_t> dependencies;
+  HIP_CHECK(hipGraphCreate(&graph, 0));
+
+  // Adding MemcpyNode
+  HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, 0, A_d, A_h,
+                                    Nbytes, hipMemcpyHostToDevice));
+  dependencies.push_back(memcpyH2D_A);
+
+  HIP_CHECK(hipGraphAddMemcpyNodeToSymbol(&memcpyToSymbolNode, graph,
+          dependencies.data(),
+          dependencies.size(),
+          HIP_SYMBOL(globalIn),
+          A_d, Nbytes, 0,
+          hipMemcpyDeviceToDevice));
+  dependencies.clear();
+  dependencies.push_back(memcpyToSymbolNode);
+
+
+  HIP_CHECK(hipGraphAddMemcpyNodeFromSymbol(&memcpyFromSymbolNode, graph,
+                                            dependencies.data(),
+                                            dependencies.size(),
+                                            B_d,
+                                            HIP_SYMBOL(globalIn),
+                                            Nbytes, 0,
+                                            hipMemcpyDeviceToDevice));
+  dependencies.clear();
+  dependencies.push_back(memcpyFromSymbolNode);
+
+  // Adding Kernel node
+  void* kernelArgs1[] = {&B_d};
+  kernelNodeParams.func =
+                       reinterpret_cast<void *>(MemcpyFromSymbolKernel);
+  kernelNodeParams.gridDim = dim3(blocks);
+  kernelNodeParams.blockDim = dim3(threadsPerBlock);
+  kernelNodeParams.sharedMemBytes = 0;
+  kernelNodeParams.kernelParams = reinterpret_cast<void**>(kernelArgs1);
+  kernelNodeParams.extra = nullptr;
+  HIP_CHECK(hipGraphAddKernelNode(&memcpyfromsymbolkernel, graph,
+                                  dependencies.data(), dependencies.size(),
+                                  &kernelNodeParams));
+  dependencies.clear();
+  dependencies.push_back(memcpyfromsymbolkernel);
+
+  // Adding MemcpyNode
+  HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_B, graph, dependencies.data(),
+                                    dependencies.size(), B_h, B_d,
+                                    Nbytes, hipMemcpyDeviceToHost));
+
+  // Instantiate and launch the graph
+  HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
+  HIP_CHECK(hipGraphLaunch(graphExec, 0));
+
+  // Validating the result
+  for (int i = 0; i < SIZE; i++) {
+    if (B_h[i] != A_h[i]) {
+       WARN("Validation failed B_h[i] " << B_h[i] << "A_h[i] " << A_h[i]);
+       REQUIRE(false);
+    }
+  }
+
+  HipTest::freeArrays<int>(A_d, B_d, nullptr,
+                           A_h, B_h, nullptr, false);
+  HIP_CHECK(hipGraphExecDestroy(graphExec));
+  HIP_CHECK(hipGraphDestroy(graph));
+}

tests/catch/unit/graph/hipGraphAddMemcpyNodeToSymbol.cc

@@ -0,0 +1,400 @@
+/*
+Copyright (c) 2021 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, INCLUDING 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 ANY 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.
+*/
+
+/**
+Testcase Scenarios of hipGraphAddMemcpyNodeToSymbol API:
+
+Functional :
+
+1. Allocate global symbol memory, add the MemcpyNodeToSymbol
+   node to the graph and verify  for different memory kinds
+2. Allocate const memory  add the MemcpyNodeToSymbol node to
+   the graph and verify  for different memory kinds
+3. Allocate global symbol memory and  device memory in GPU-0
+   and perform MemcpyToSymbol from peer GPU by adding it to the graph node.
+4. Allocate const symbol memory and  device memory in GPU-0
+   and perform MemcpyToSymbol from peer GPU by adding it to the graph node.
+5. Allocate global memory, Add MemcpyToSymbolNode,KernelNode and memcpynode and validating
+   the behaviour
+
+Negative :
+
+1) Pass nullptr to graph node
+2) Pass nullptr to graph
+3) Pass nullptr to dependencies
+4) Pass invalid numDependencies
+5) Pass nullptr to dst
+6) Pass nullptr to symbol
+7) Pass invalid count 
+8) Pass offset+count greater than allocated size
+9) Pass unintialized graph
+*/
+
+#include <hip_test_common.hh>
+#include <hip_test_checkers.hh>
+#include <limits>
+#define SIZE 256
+
+__device__ int globalIn[SIZE];
+__device__ __constant__ int globalConst[SIZE];
+
+__global__ void MemcpyToSymbolKernel(int* B_d) {
+  for (int i = 0 ; i < SIZE; i++) {
+      B_d[i] = globalIn[i];
+  }
+}
+
+/* This testcase verifies negative scenarios of
+   hipGraphAddMemcpyNodeToSymbol API */
+TEST_CASE("Unit_hipGraphAddMemcpyNodeToSymbol_Negative") {
+  constexpr size_t Nbytes = SIZE * sizeof(int);
+  int *A_d{nullptr};
+  int *A_h{nullptr}, *B_h{nullptr};
+  HipTest::initArrays<int>(&A_d, nullptr, nullptr,
+                           &A_h, &B_h, nullptr, SIZE, false);
+
+  hipGraph_t graph;
+  hipGraphNode_t memcpyH2D_A;
+  std::vector<hipGraphNode_t> dependencies;
+  HIP_CHECK(hipGraphCreate(&graph, 0));
+
+  // Adding MemcpyNode
+  HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, 0, A_d, A_h,
+                                    Nbytes, hipMemcpyHostToDevice));
+  dependencies.push_back(memcpyH2D_A);
+#if HT_NVIDIA
+  hipGraphNode_t memcpyToSymbolNode;
+  SECTION("Passing nullptr to graph") {
+    REQUIRE(hipGraphAddMemcpyNodeToSymbol(&memcpyToSymbolNode, nullptr,
+                                          dependencies.data(),
+                                          dependencies.size(),
+                                          HIP_SYMBOL(globalIn),
+                                          A_h, Nbytes, 0,
+                                          hipMemcpyDeviceToDevice)
+                                          == hipErrorInvalidValue);
+  }
+
+  SECTION("Passing nullptr to graph node") {
+    REQUIRE(hipGraphAddMemcpyNodeToSymbol(nullptr, graph,
+                                          dependencies.data(),
+                                          dependencies.size(),
+                                          HIP_SYMBOL(globalIn),
+                                          A_d, Nbytes, 0,
+                                          hipMemcpyDeviceToDevice)
+                                          == hipErrorInvalidValue);
+  }
+
+  SECTION("Passing size > 1 and dependencies as nullptr") {
+    REQUIRE(hipGraphAddMemcpyNodeToSymbol(&memcpyToSymbolNode, graph,
+                                          nullptr,
+                                          1,
+                                          HIP_SYMBOL(globalIn),
+                                          A_d, Nbytes, 0,
+                                          hipMemcpyDeviceToDevice)
+                                          == hipErrorInvalidValue);
+  }
+
+  SECTION("Passing invalid dependencies size") {
+    REQUIRE(hipGraphAddMemcpyNodeToSymbol(&memcpyToSymbolNode, graph,
+                                          dependencies.data(),
+                                          10,
+                                          HIP_SYMBOL(globalIn),
+                                          A_d, Nbytes, 0,
+                                          hipMemcpyDeviceToDevice)
+                                          == hipErrorInvalidValue);
+  }
+
+  SECTION("Passing nullptr to dst") {
+    REQUIRE(hipGraphAddMemcpyNodeToSymbol(&memcpyToSymbolNode, graph,
+                                          dependencies.data(),
+                                          dependencies.size(),
+                                          nullptr,
+                                          A_d, Nbytes, 0,
+                                          hipMemcpyDeviceToDevice)
+                                          == hipErrorInvalidSymbol);
+  }
+
+  SECTION("Passing nullptr to source") {
+    REQUIRE(hipGraphAddMemcpyNodeToSymbol(&memcpyToSymbolNode, graph,
+                                          dependencies.data(),
+                                          dependencies.size(),
+                                          HIP_SYMBOL(globalIn),
+                                          nullptr, Nbytes, 0,
+                                          hipMemcpyDeviceToDevice)
+                                          == hipErrorInvalidValue);
+  }
+
+  SECTION("Passing offset+size > max size") {
+    REQUIRE(hipGraphAddMemcpyNodeToSymbol(&memcpyToSymbolNode, graph,
+                                          dependencies.data(),
+                                          dependencies.size(),
+                                          HIP_SYMBOL(globalIn),
+                                          A_d, Nbytes, 10,
+                                          hipMemcpyDeviceToDevice)
+                                          == hipErrorInvalidValue);
+  }
+
+  SECTION("Passing Max count") {
+  REQUIRE(hipGraphAddMemcpyNodeToSymbol(&memcpyToSymbolNode, graph,
+                                        dependencies.data(),
+                                        dependencies.size(),
+                                        HIP_SYMBOL(globalIn),
+                                        A_d,
+                                        std::numeric_limits<int>::max(), 0,
+                                        hipMemcpyDeviceToDevice)
+                                        == hipErrorInvalidValue);
+  }
+
+  SECTION("Pass Unintialized graph") {
+  hipGraph_t unint_graph;
+  REQUIRE(hipGraphAddMemcpyNodeToSymbol(&memcpyToSymbolNode, unint_graph,
+                                        dependencies.data(),
+                                        dependencies.size(),
+                                        HIP_SYMBOL(globalIn),
+                                        A_d,
+                                        Nbytes, 0,
+                                        hipMemcpyDeviceToDevice)
+                                        == hipErrorInvalidValue);
+  }
+#endif
+
+  HipTest::freeArrays<int>(A_d, nullptr, nullptr,
+                      A_h, B_h, nullptr, false);
+  HIP_CHECK(hipGraphDestroy(graph));
+}
+/*
+This function is used to verify the following scenarios
+1. Create global variable, allocate Memory in GPU-0 and create dependency graph of
+   hipGraphAddMemcpyNodeToSymbol API in GPU-1 and validate the result
+2. Allocate global memory, Create dependency graph and validate the result on GPU-0
+3. Allocate global const memory, Create dependency graph and validate the result on GPU-0
+4. Create global const variable, allocate Memory in GPU-0 and create dependency graph of
+   hipGraphAddMemcpyNodeToSymbol API in GPU-1 and validate the result
+*/
+void hipGraphAddMemcpyNodeToSymbol_GlobalMemory(bool device_ctxchg = false,
+                                                bool const_device_var = false) {
+  constexpr size_t Nbytes = SIZE * sizeof(int);
+  int *A_d{nullptr};
+  int *A_h{nullptr}, *B_h{nullptr};
+  HipTest::initArrays<int>(&A_d, nullptr, nullptr,
+                           &A_h, &B_h, nullptr, SIZE, false);
+
+  hipGraph_t graph;
+  hipGraphExec_t graphExec;
+  hipGraphNode_t memcpyToSymbolNode, memcpyFromSymbolNode, memcpyH2D_A;
+  std::vector<hipGraphNode_t> dependencies;
+  HIP_CHECK(hipGraphCreate(&graph, 0));
+
+  if (device_ctxchg) {
+    HIP_CHECK(hipSetDevice(1));
+    HIP_CHECK(hipDeviceEnablePeerAccess(0, 0));
+  }
+  // Adding MemcpyNode
+  HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, 0, A_d, A_h,
+                                    Nbytes, hipMemcpyHostToDevice));
+  dependencies.push_back(memcpyH2D_A);
+
+  // Adding MemcpyNodeToSymbol
+
+  if (const_device_var) {
+    HIP_CHECK(hipGraphAddMemcpyNodeToSymbol(&memcpyToSymbolNode, graph,
+                                            dependencies.data(),
+                                            dependencies.size(),
+                                            HIP_SYMBOL(globalConst),
+                                            A_d, Nbytes, 0,
+                                            hipMemcpyDeviceToDevice));
+  } else {
+    HIP_CHECK(hipGraphAddMemcpyNodeToSymbol(&memcpyToSymbolNode, graph,
+                                            dependencies.data(),
+                                            dependencies.size(),
+                                            HIP_SYMBOL(globalIn),
+                                            A_d, Nbytes, 0,
+                                            hipMemcpyDeviceToDevice));
+  }
+  dependencies.clear();
+  dependencies.push_back(memcpyToSymbolNode);
+
+  // Adding MemcpyNodeFromSymbol
+  if (const_device_var) {
+    HIP_CHECK(hipGraphAddMemcpyNodeFromSymbol(&memcpyFromSymbolNode, graph,
+          dependencies.data(),
+          dependencies.size(),
+          B_h,
+          HIP_SYMBOL(globalConst),
+          Nbytes, 0, hipMemcpyDeviceToHost));
+  } else {
+    HIP_CHECK(hipGraphAddMemcpyNodeFromSymbol(&memcpyFromSymbolNode, graph,
+          dependencies.data(),
+          dependencies.size(),
+          B_h,
+          HIP_SYMBOL(globalIn),
+          Nbytes, 0, hipMemcpyDeviceToHost));
+  }
+
+  // Instantiate and launch the graph
+  HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
+  HIP_CHECK(hipGraphLaunch(graphExec, 0));
+
+  // Validating the result
+  for (int i = 0; i < SIZE; i++) {
+    if (B_h[i] != A_h[i]) {
+       WARN("Validation failed B_h[i] " << B_h[i] << "A_h[i] " << A_h[i]);
+       REQUIRE(false);
+    }
+  }
+
+  HipTest::freeArrays<int>(A_d, nullptr, nullptr,
+                      A_h, B_h, nullptr, false);
+  HIP_CHECK(hipGraphExecDestroy(graphExec));
+  HIP_CHECK(hipGraphDestroy(graph));
+}
+/*
+This testcase verifies allocating global symbol memory,
+add the MemcpyNodeToSymbol node to the graph and
+erifying the result
+*/
+TEST_CASE("Unit_hipGraphAddMemcpyNodeToSymbol_GlobalMemory") {
+  hipGraphAddMemcpyNodeToSymbol_GlobalMemory(false, false);
+}
+
+/*
+This testcase verifies allocating global const symbol memory,
+add the MemcpyNodeToSymbol node to the graph and
+verifying the result
+*/
+TEST_CASE("Unit_hipGraphAddMemcpyNodeToSymbol_GlobalConstMemory") {
+  hipGraphAddMemcpyNodeToSymbol_GlobalMemory(false, true);
+}
+
+#if HT_NVIDIA
+/*
+This testcase verifies allocating global symbol memory and device variables
+in GPU-0 and add the MemcpyNodeToSymbol node to the graph and
+verifying the result in GPU-1
+*/
+TEST_CASE("Unit_hipGraphAddMemcpyNodeToSymbol_GlobalMemoryPeerDevice") {
+  int numDevices = 0;
+  int canAccessPeer = 0;
+  HIP_CHECK(hipGetDeviceCount(&numDevices));
+  if (numDevices > 1) {
+    hipDeviceCanAccessPeer(&canAccessPeer, 0, 1);
+    if (canAccessPeer) {
+      hipGraphAddMemcpyNodeToSymbol_GlobalMemory(true, false);
+    } else {
+      SUCCEED("Machine does not seem to have P2P");
+    }
+  } else {
+    SUCCEED("skipped the testcase as no of devices is less than 2");
+  }
+}
+/*
+This testcase verifies allocating global const symbol memory and device variables
+in GPU-0 and add the MemcpyNodeToSymbol node to the graph and
+verifying the result in GPU-1
+*/
+TEST_CASE("Unit_hipGraphAddMemcpyNodeToSymbol_GlobalConstMemoryPeerDevice") {
+  int numDevices = 0;
+  int canAccessPeer = 0;
+  HIP_CHECK(hipGetDeviceCount(&numDevices));
+  if (numDevices > 1) {
+    hipDeviceCanAccessPeer(&canAccessPeer, 0, 1);
+    if (canAccessPeer) {
+      hipGraphAddMemcpyNodeToSymbol_GlobalMemory(true, true);
+    } else {
+      SUCCEED("Machine does not seem to have P2P");
+    }
+  } else {
+    SUCCEED("skipped the testcase as no of devices is less than 2");
+  }
+}
+#endif
+/*
+This testcaser verifies allocating global memory,
+Add MemcpyToSymbolNode,KernelNode and memcpynode and validating
+the behaviour
+*/
+TEST_CASE("Unit_hipGraphAddMemcpyNodeToSymbol_MemcpyToSymbolNodeWithKernel") {
+  constexpr size_t Nbytes = SIZE * sizeof(int);
+  constexpr auto blocksPerCU = 6;  // to hide latency
+  constexpr auto threadsPerBlock = 256;
+  unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, SIZE);
+  hipGraphNode_t memcpytosymbolkernel, memcpyD2H_B;
+  hipKernelNodeParams kernelNodeParams{};
+  int *A_d{nullptr}, *B_d{nullptr};
+  int *A_h{nullptr}, *B_h{nullptr};
+  HipTest::initArrays<int>(&A_d, &B_d, nullptr,
+                           &A_h, &B_h, nullptr, SIZE, false);
+
+  hipGraph_t graph;
+  hipGraphExec_t graphExec;
+  hipGraphNode_t memcpyToSymbolNode, memcpyH2D_A;
+  std::vector<hipGraphNode_t> dependencies;
+  HIP_CHECK(hipGraphCreate(&graph, 0));
+
+  // Adding MemcpyNode
+  HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, 0, A_d, A_h,
+                                    Nbytes, hipMemcpyHostToDevice));
+  dependencies.push_back(memcpyH2D_A);
+
+  HIP_CHECK(hipGraphAddMemcpyNodeToSymbol(&memcpyToSymbolNode, graph,
+                                            dependencies.data(),
+                                            dependencies.size(),
+                                            HIP_SYMBOL(globalIn),
+                                            A_d, Nbytes, 0,
+                                            hipMemcpyDeviceToDevice));
+  dependencies.clear();
+  dependencies.push_back(memcpyToSymbolNode);
+
+  // Adding Kernel node
+  void* kernelArgs1[] = {&B_d};
+  kernelNodeParams.func =
+                       reinterpret_cast<void *>(MemcpyToSymbolKernel);
+  kernelNodeParams.gridDim = dim3(blocks);
+  kernelNodeParams.blockDim = dim3(threadsPerBlock);
+  kernelNodeParams.sharedMemBytes = 0;
+  kernelNodeParams.kernelParams = reinterpret_cast<void**>(kernelArgs1);
+  kernelNodeParams.extra = nullptr;
+  HIP_CHECK(hipGraphAddKernelNode(&memcpytosymbolkernel, graph,
+                                  dependencies.data(), dependencies.size(),
+                                  &kernelNodeParams));
+  dependencies.clear();
+  dependencies.push_back(memcpytosymbolkernel);
+
+  // Adding MemcpyNode
+  HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_B, graph, dependencies.data(),
+                                    dependencies.size(), B_h, B_d,
+                                    Nbytes, hipMemcpyDeviceToHost));
+
+  // Instantiate and launch the graph
+  HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
+  HIP_CHECK(hipGraphLaunch(graphExec, 0));
+
+  // Validating the result
+  for (int i = 0; i < SIZE; i++) {
+    if (B_h[i] != A_h[i]) {
+       WARN("Validation failed B_h[i] " << B_h[i] << "A_h[i] " << A_h[i]);
+       REQUIRE(false);
+    }
+  }
+
+  HipTest::freeArrays<int>(A_d, B_d, nullptr,
+                           A_h, B_h, nullptr, false);
+  HIP_CHECK(hipGraphExecDestroy(graphExec));
+  HIP_CHECK(hipGraphDestroy(graph));
+}

tests/catch/unit/graph/hipGraphChildGraphNodeGetGraph.cc

@@ -0,0 +1,163 @@
+/*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, INCLUDING 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 ANY 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.
+*/
+
+/**
+Testcase Scenarios of hipGraphChildGraphNodeGetGraph API:
+
+Functional Scenarios:
+1. Get the child graph node from the original graph and execute it
+
+Negative Scenarios:
+1. Pass nullptr to graph
+2. Pass nullptr to graphnode
+3. Pass uninitialized graph node
+4. Pass orginial graph node instead of child graph node
+**/
+
+#include <hip_test_common.hh>
+#include <hip_test_checkers.hh>
+#include <hip_test_kernels.hh>
+
+/*
+This testcase verifies the following scenario
+Create graph, add multiple child nodes and gets the
+graph of one of the child nodes using hipGraphChildGraphNodeGetGraph API
+executes it and validates the results
+*/
+TEST_CASE("Unit_hipGraphChildGraphNodeGetGraph_Functional") {
+  constexpr size_t N = 1024;
+  constexpr size_t Nbytes = N * sizeof(int);
+  constexpr auto blocksPerCU = 6;  // to hide latency
+  size_t NElem{N};
+  constexpr auto threadsPerBlock = 256;
+  hipGraph_t graph, childgraph1, childgraph2;
+  hipGraphExec_t graphExec;
+  hipKernelNodeParams kernelNodeParams{};
+  hipGraphNode_t kernel_vecAdd;
+  int *A_d{nullptr}, *B_d{nullptr}, *C_d{nullptr};
+  int *A_h{nullptr}, *B_h{nullptr}, *C_h{nullptr};
+  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));
+  hipGraphNode_t memcpyH2D_A, memcpyH2D_B, childGraphNode1,
+                 childGraphNode2, memcpyD2H_C;
+  hipStream_t streamForGraph;
+  HIP_CHECK(hipStreamCreate(&streamForGraph));
+  HIP_CHECK(hipGraphCreate(&childgraph1, 0));
+  HIP_CHECK(hipGraphCreate(&childgraph2, 0));
+  HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, childgraph1, nullptr,
+                                    0, A_d, A_h,
+                                    Nbytes, hipMemcpyHostToDevice));
+  HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_B, childgraph2, nullptr,
+                                    0, B_d, B_h,
+                                    Nbytes, hipMemcpyHostToDevice));
+
+  HIP_CHECK(hipGraphAddChildGraphNode(&childGraphNode1, graph,
+                                      nullptr, 0, childgraph1));
+  HIP_CHECK(hipGraphAddChildGraphNode(&childGraphNode2, graph,
+                                      nullptr, 0, childgraph2));
+  HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_C, graph, nullptr,
+                                    0, C_h, C_d,
+                                    Nbytes, hipMemcpyDeviceToHost));
+
+  void* kernelArgs2[] = {&A_d, &B_d, &C_d, reinterpret_cast<void *>(&NElem)};
+  kernelNodeParams.func = reinterpret_cast<void *>(HipTest::vectorADD<int>);
+  kernelNodeParams.gridDim = dim3(blocks);
+  kernelNodeParams.blockDim = dim3(threadsPerBlock);
+  kernelNodeParams.sharedMemBytes = 0;
+  kernelNodeParams.kernelParams = reinterpret_cast<void**>(kernelArgs2);
+  kernelNodeParams.extra = nullptr;
+  HIP_CHECK(hipGraphAddKernelNode(&kernel_vecAdd, graph, nullptr, 0,
+                                                        &kernelNodeParams));
+  HIP_CHECK(hipGraphAddDependencies(graph, &childGraphNode1,
+                                    &childGraphNode2, 1));
+  HIP_CHECK(hipGraphAddDependencies(graph, &childGraphNode2,
+                                    &kernel_vecAdd, 1));
+  HIP_CHECK(hipGraphAddDependencies(graph, &kernel_vecAdd, &memcpyD2H_C, 1));
+  hipGraph_t Getgraph;
+  HIP_CHECK(hipGraphChildGraphNodeGetGraph(childGraphNode1, &Getgraph));
+  // Instantiate and launch the child graph
+  HIP_CHECK(hipGraphInstantiate(&graphExec, Getgraph, nullptr, nullptr, 0));
+  HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
+  HIP_CHECK(hipStreamSynchronize(streamForGraph));
+
+
+  // Verify child graph execution result
+  HIP_CHECK(hipMemcpy(C_h, A_d, Nbytes, hipMemcpyDeviceToHost));
+  for (size_t i = 0; i < N; i++) {
+    if (A_h[i] != C_h[i]) {
+      INFO("Validation failed " << A_h[i] << C_h[i]);
+      REQUIRE(false);
+    }
+  }
+  HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
+  HIP_CHECK(hipGraphExecDestroy(graphExec));
+  HIP_CHECK(hipGraphDestroy(childgraph2));
+  HIP_CHECK(hipGraphDestroy(childgraph1));
+  HIP_CHECK(hipGraphDestroy(graph));
+  HIP_CHECK(hipStreamDestroy(streamForGraph));
+}
+
+/*
+This testcase verifies the negative scenarios
+of hipGraphChildGraphNodeGetGraph API
+*/
+TEST_CASE("Unit_hipGraphChildGraphNodeGetGraph_Negative") {
+  constexpr size_t N = 1024;
+  constexpr size_t Nbytes = N * sizeof(int);
+  hipGraph_t graph, childgraph1;
+  int *A_d{nullptr}, *B_d{nullptr}, *C_d{nullptr};
+  int *A_h{nullptr}, *B_h{nullptr}, *C_h{nullptr};
+  HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
+
+  HIP_CHECK(hipGraphCreate(&graph, 0));
+  hipGraphNode_t memcpyH2D_A, childGraphNode1;
+  HIP_CHECK(hipGraphCreate(&childgraph1, 0));
+  HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, childgraph1, nullptr,
+                                    0, A_d, A_h,
+                                    Nbytes, hipMemcpyHostToDevice));
+
+  HIP_CHECK(hipGraphAddChildGraphNode(&childGraphNode1, graph,
+                                      nullptr, 0, childgraph1));
+
+  hipGraph_t Getgraph;
+  SECTION("nullptr to child node") {
+    REQUIRE((hipGraphChildGraphNodeGetGraph(nullptr, &Getgraph))
+                                            == hipErrorInvalidValue);
+  }
+#if HT_NVIDIA
+  SECTION("nullptr to graph") {
+    REQUIRE((hipGraphChildGraphNodeGetGraph(childGraphNode1, nullptr))
+                                            == hipErrorInvalidValue);
+  }
+
+  SECTION("Passing parent instead of child graph node") {
+    REQUIRE((hipGraphChildGraphNodeGetGraph(memcpyH2D_A, &Getgraph))
+                                            == hipErrorInvalidValue);
+  }
+
+  SECTION("Passing unintialized node") {
+    hipGraphNode_t unint_node;
+    REQUIRE((hipGraphChildGraphNodeGetGraph(unint_node, &Getgraph))
+                                            == hipErrorInvalidValue);
+  }
+#endif
+  HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
+  HIP_CHECK(hipGraphDestroy(childgraph1));
+}

tests/catch/unit/graph/hipGraphClone.cc

@@ -0,0 +1,319 @@
+/*
+Copyright (c) 2021 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, INCLUDING 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 ANY 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.
+*/
+
+/*
+Testcase Scenarios of hipGraphClone API:
+
+Negative:
+
+1. Pass nullptr to cloned graph
+2. pass nullptr to original graph
+
+Functional:
+
+1. Clone the graph,Instantiate and execute the cloned graph
+2. Clone the graph and modify the original graph and ensure that the
+   cloned graph is not modified
+3. Create graph on one GPU device and clone it from peer GPU device
+4. Create graph in one thread and clone it from multiple threads.
+*/
+
+#include <hip_test_common.hh>
+#include <hip_test_checkers.hh>
+#include <hip_test_kernels.hh>
+
+#define NUM_THREADS 10
+
+/* This test covers the negative scenarios of
+   hipGraphClone API */
+
+TEST_CASE("Unit_hipGraphClone_Negative") {
+  SECTION("Passing nullptr to Cloned graph") {
+    hipGraph_t graph;
+    HIP_CHECK(hipGraphCreate(&graph, 0));
+    REQUIRE(hipGraphClone(nullptr, graph) == hipErrorInvalidValue);
+    HIP_CHECK(hipGraphDestroy(graph));
+  }
+
+  SECTION("Passing nullptr to original graph") {
+    hipGraph_t clonedGraph;
+    REQUIRE(hipGraphClone(&clonedGraph, nullptr) == hipErrorInvalidValue);
+  }
+}
+/*
+This function creates the graph with dependencies
+then performs device context change and clones the cloned graph
+Executes the cloned graph and validates the result
+*/
+void hipGraphClone_DeviceContextChange() {
+  constexpr size_t N = 1024;
+  constexpr size_t Nbytes = N * sizeof(int);
+  hipGraph_t graph, clonedgraph;
+  hipGraphExec_t graphExec;
+  hipStream_t streamForGraph;
+  hipGraphNode_t memcpyH2D_A, memcpyD2H_A;
+  int *A_d{nullptr}, *A_h{nullptr}, *B_h{nullptr};
+  HipTest::initArrays<int>(&A_d, nullptr, nullptr,
+                      &A_h, &B_h, nullptr, N, false);
+  HIP_CHECK(hipGraphCreate(&graph, 0));
+  HIP_CHECK(hipStreamCreate(&streamForGraph));
+  HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, 0, A_d, A_h,
+                                    Nbytes, hipMemcpyHostToDevice));
+  HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_A, graph, nullptr, 0, B_h, A_d,
+                                    Nbytes, hipMemcpyDeviceToHost));
+  HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_A, &memcpyD2H_A, 1));
+  HIP_CHECK(hipSetDevice(1));
+  HIP_CHECK(hipGraphClone(&clonedgraph, graph));
+  // Instantiate and launch the original graph
+  HIP_CHECK(hipGraphInstantiate(&graphExec, clonedgraph, nullptr, nullptr, 0));
+  HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
+  HIP_CHECK(hipStreamSynchronize(streamForGraph));
+
+  for (size_t i = 0; i < N; i++) {
+    if (A_h[i] != B_h[i]) {
+      INFO("Validation failed A_h[i] " << A_h[i] << " B_h[i] " << B_h[i]);
+      REQUIRE(false);
+    }
+  }
+  HipTest::freeArrays<int>(A_d, nullptr, nullptr, A_h, B_h, nullptr, false);
+  HIP_CHECK(hipGraphExecDestroy(graphExec));
+  HIP_CHECK(hipGraphDestroy(graph));
+  HIP_CHECK(hipGraphDestroy(clonedgraph));
+  HIP_CHECK(hipStreamDestroy(streamForGraph));
+}
+/*
+This function does the following
+1. Creates the graph with multiple dependencies
+   clones the graph and validates the result.
+2. Creates the graph, clones the graph and modifies
+   the existing graph and execute the cloned graph
+   to ensure that cloned graph is not modified
+*/
+void hipGraphClone_Func(bool ModifyOrigGraph = false) {
+  constexpr size_t N = 1024;
+  constexpr size_t Nbytes = N * sizeof(int);
+  constexpr auto blocksPerCU = 6;  // to hide latency
+  constexpr auto threadsPerBlock = 256;
+  hipGraph_t graph, clonedgraph;
+  hipGraphNode_t memset_A, memset_B, memsetKer_C;
+  hipGraphNode_t memcpyH2D_A, memcpyH2D_B, memcpyD2H_C, memcpyD2D_C,
+                 memcpyD2H_C_new;
+  hipGraphNode_t kernel_vecAdd;
+  hipKernelNodeParams kernelNodeParams{};
+  hipStream_t streamForGraph;
+  int *A_d, *B_d, *C_d;
+  int *A_h, *B_h, *C_h;
+  hipGraphExec_t graphExec;
+  hipMemsetParams memsetParams{};
+  int memsetVal{};
+  size_t NElem{N};
+
+  HIP_CHECK(hipStreamCreate(&streamForGraph));
+  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));
+
+  memset(&memsetParams, 0, sizeof(memsetParams));
+  memsetParams.dst = reinterpret_cast<void*>(A_d);
+  memsetParams.value = 0;
+  memsetParams.pitch = 0;
+  memsetParams.elementSize = sizeof(char);
+  memsetParams.width = Nbytes;
+  memsetParams.height = 1;
+  HIP_CHECK(hipGraphAddMemsetNode(&memset_A, graph, nullptr, 0,
+                                                              &memsetParams));
+
+  memset(&memsetParams, 0, sizeof(memsetParams));
+  memsetParams.dst = reinterpret_cast<void*>(B_d);
+  memsetParams.value = 0;
+  memsetParams.pitch = 0;
+  memsetParams.elementSize = sizeof(char);
+  memsetParams.width = Nbytes;
+  memsetParams.height = 1;
+  HIP_CHECK(hipGraphAddMemsetNode(&memset_B, graph, nullptr, 0,
+                                                              &memsetParams));
+
+  void* kernelArgs1[] = {&C_d, &memsetVal, reinterpret_cast<void *>(&NElem)};
+  kernelNodeParams.func =
+                       reinterpret_cast<void *>(HipTest::memsetReverse<int>);
+  kernelNodeParams.gridDim = dim3(blocks);
+  kernelNodeParams.blockDim = dim3(threadsPerBlock);
+  kernelNodeParams.sharedMemBytes = 0;
+  kernelNodeParams.kernelParams = reinterpret_cast<void**>(kernelArgs1);
+  kernelNodeParams.extra = nullptr;
+  HIP_CHECK(hipGraphAddKernelNode(&memsetKer_C, graph, nullptr, 0,
+                                                        &kernelNodeParams));
+
+  HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, 0, A_d, A_h,
+                                   Nbytes, hipMemcpyHostToDevice));
+
+  HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_B, graph, nullptr, 0, B_d, B_h,
+                                   Nbytes, hipMemcpyHostToDevice));
+
+  HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_C, graph, nullptr, 0, C_h, C_d,
+                                   Nbytes, hipMemcpyDeviceToHost));
+
+  void* kernelArgs2[] = {&A_d, &B_d, &C_d, reinterpret_cast<void *>(&NElem)};
+  kernelNodeParams.func = reinterpret_cast<void *>(HipTest::vectorADD<int>);
+  kernelNodeParams.gridDim = dim3(blocks);
+  kernelNodeParams.blockDim = dim3(threadsPerBlock);
+  kernelNodeParams.sharedMemBytes = 0;
+  kernelNodeParams.kernelParams = reinterpret_cast<void**>(kernelArgs2);
+  kernelNodeParams.extra = nullptr;
+  HIP_CHECK(hipGraphAddKernelNode(&kernel_vecAdd, graph, nullptr, 0,
+                                                        &kernelNodeParams));
+
+  // Create dependencies
+  HIP_CHECK(hipGraphAddDependencies(graph, &memset_A, &memcpyH2D_A, 1));
+  HIP_CHECK(hipGraphAddDependencies(graph, &memset_B, &memcpyH2D_B, 1));
+  HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_A, &kernel_vecAdd, 1));
+  HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_B, &kernel_vecAdd, 1));
+  HIP_CHECK(hipGraphAddDependencies(graph, &memsetKer_C, &kernel_vecAdd, 1));
+  HIP_CHECK(hipGraphAddDependencies(graph, &kernel_vecAdd, &memcpyD2H_C, 1));
+
+  HIP_CHECK(hipGraphClone(&clonedgraph, graph));
+
+  if (ModifyOrigGraph) {
+    // Modify Original graph by adding new dependency
+    HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2D_C, graph, nullptr, 0,
+                                      C_d, B_d,
+                                      Nbytes, hipMemcpyDeviceToHost));
+
+    HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_C_new, graph, nullptr, 0,
+                                      C_h, C_d,
+                                      Nbytes, hipMemcpyDeviceToHost));
+
+    HIP_CHECK(hipGraphAddDependencies(graph, &kernel_vecAdd, &memcpyD2D_C, 1));
+    HIP_CHECK(hipGraphAddDependencies(graph, &memcpyD2D_C,
+                                      &memcpyD2H_C_new, 1));
+
+    // Instantiate and launch the original graph
+    HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
+    HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
+    HIP_CHECK(hipStreamSynchronize(streamForGraph));
+
+    for (size_t i= 0; i < NElem; i++) {
+      if (C_h[i] != B_h[i]) {
+         INFO("Validation failed C_h is " << C_h[i] <<
+               "B_h is " << B_h[i]);
+         REQUIRE(false);
+      }
+    }
+  }
+
+  // Instantiate and launch the cloned graph
+  HIP_CHECK(hipGraphInstantiate(&graphExec, clonedgraph, nullptr, nullptr, 0));
+  HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
+  HIP_CHECK(hipStreamSynchronize(streamForGraph));
+
+  // 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(hipGraphDestroy(clonedgraph));
+  HIP_CHECK(hipStreamDestroy(streamForGraph));
+}
+
+/*
+This testcase verifies following scenarios
+1. Clones the graph and verify the result
+2. Clones the graph, Modify the original graph and
+   validate the result of the cloned graph
+3. Device context change for cloned graph
+*/
+TEST_CASE("Unit_hipGraphClone_Functional") {
+  SECTION("hipGraphClone Basic Functionality") {
+    hipGraphClone_Func();
+  }
+  SECTION("hipGraphClone Modify Original graph") {
+    hipGraphClone_Func(true);
+  }
+
+  SECTION("hipGraphClone Device context change") {
+    int numDevices = 0;
+    int canAccessPeer = 0;
+    HIP_CHECK(hipGetDeviceCount(&numDevices));
+    if (numDevices > 1) {
+      HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1));
+      if (canAccessPeer) {
+        hipGraphClone_DeviceContextChange();
+      } else {
+        SUCCEED("Machine does not seem to have P2P");
+      }
+    } else {
+      SUCCEED("skipped the testcase as no of devices is less than 2");
+    }
+  }
+}
+
+/*
+This testcase creates the graph with dependencies
+then creates multiple threads and clones the graph
+in each thread and executes the cloned graph
+hipGraphClone is failing in CUDA in multi threaded
+scenario so excluded for nvidia
+*/
+#if HT_AMD
+TEST_CASE("Unit_hipGraphClone_MultiThreaded") {
+  constexpr size_t N = 1024;
+  constexpr size_t Nbytes = N * sizeof(int);
+  hipGraph_t graph;
+  hipGraphNode_t memcpyH2D_A, memcpyD2H_A;
+  int *A_d{nullptr}, *A_h{nullptr}, *B_h{nullptr};
+  HipTest::initArrays<int>(&A_d, nullptr, nullptr,
+                      &A_h, &B_h, nullptr, N, false);
+  HIP_CHECK(hipGraphCreate(&graph, 0));
+  HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, 0, A_d, A_h,
+                                    Nbytes, hipMemcpyHostToDevice));
+  HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_A, graph, nullptr, 0, B_h, A_d,
+                                    Nbytes, hipMemcpyDeviceToHost));
+  HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_A, &memcpyD2H_A, 1));
+  std::vector<std::thread> threads;
+  auto lambdaFunc = [&](){
+    hipGraph_t clonedgraph;
+    hipGraphExec_t graphExec;
+    HIP_CHECK(hipGraphClone(&clonedgraph, graph));
+    // Instantiate and launch the cloned graph
+    HIP_CHECK(hipGraphInstantiate(&graphExec, clonedgraph, nullptr,
+          nullptr, 0));
+    HIP_CHECK(hipGraphLaunch(graphExec, 0));
+
+    for (size_t i = 0; i < N; i++) {
+      if (A_h[i] != B_h[i]) {
+        INFO("Validation failed A_h[i] " << A_h[i] << " B_h[i] " << B_h[i]);
+        REQUIRE(false);
+      }
+    }
+
+    HIP_CHECK(hipGraphExecDestroy(graphExec));
+    HIP_CHECK(hipGraphDestroy(clonedgraph));
+  };
+  for (int i = 0; i < NUM_THREADS; i++) {
+    std::thread t(lambdaFunc);
+    threads.push_back(std::move(t));
+  }
+  for (auto &t : threads) {
+    t.join();
+  }
+  HipTest::freeArrays<int>(A_d, nullptr, nullptr, A_h, B_h, nullptr, false);
+  HIP_CHECK(hipGraphDestroy(graph));
+}
+#endif

tests/catch/unit/graph/hipGraphDestroyNode.cc

@@ -0,0 +1,139 @@
+/*
+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, INCLUDING 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 ANY 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.
+*/
+
+/*
+Testcase Scenarios of hipGraphDestroyNode API:
+
+Negative ::
+
+1) Pass nullptr to graph node
+
+Functional ::
+
+1) Create Node and destroy the node
+2) Create graph with dependencies and destroy one of the dependency node
+   before executing the graph.
+*/
+
+#include <hip_test_common.hh>
+#include <hip_test_checkers.hh>
+#include <hip_test_kernels.hh>
+
+
+/* This test covers the negative scenarios of
+   hipGraphDestroyNode API */
+TEST_CASE("Unit_hipGraphDestroyNode_Negative") {
+  SECTION("Passing nullptr to graph Node") {
+    REQUIRE(hipGraphDestroyNode(nullptr) == hipErrorInvalidValue);
+  }
+}
+
+/* This test covers the basic functionality of
+   hipGraphDestroyNode API where we create and destroy
+   the node
+*/
+TEST_CASE("Unit_hipGraphDestroyNode_BasicFunctionality") {
+  char *pOutBuff_d{};
+  constexpr size_t size = 1024;
+  hipGraph_t graph{};
+  hipGraphNode_t memsetNode{};
+
+  HIP_CHECK(hipMalloc(&pOutBuff_d, size));
+  hipMemsetParams memsetParams{};
+  memsetParams.dst = reinterpret_cast<void*>(pOutBuff_d);
+  memsetParams.value = 0;
+  memsetParams.pitch = 0;
+  memsetParams.elementSize = sizeof(char);
+  memsetParams.width = size * sizeof(char);
+  memsetParams.height = 1;
+  HIP_CHECK(hipGraphCreate(&graph, 0));
+  HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0,
+                                                              &memsetParams));
+  REQUIRE(hipGraphDestroyNode(memsetNode) == hipSuccess);
+  HIP_CHECK(hipFree(pOutBuff_d));
+}
+
+/*
+This testcase verifies the following scenario where
+graph is created with dependencies and one of the dependency is
+destroyed before execute the graph
+*/
+TEST_CASE("Unit_hipGraphDestroyNode_DestroyDependencyNode") {
+  constexpr size_t N = 1024;
+  constexpr size_t Nbytes = N * sizeof(int);
+  constexpr auto blocksPerCU = 6;  // to hide latency
+  constexpr auto threadsPerBlock = 256;
+  hipGraph_t graph;
+  hipGraphNode_t memcpyH2D_A, memcpyH2D_B, memcpyH2D_B2Copies, memcpyD2H_C;
+  hipGraphNode_t kernel_vecAdd;
+  hipKernelNodeParams kernelNodeParams{};
+  int *A_d, *B_d, *C_d;
+  int *A_h, *B_h, *C_h;
+  hipGraphExec_t graphExec;
+  size_t NElem{N};
+  hipStream_t streamForGraph;
+
+  HIP_CHECK(hipStreamCreate(&streamForGraph));
+  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(&memcpyH2D_A, graph, nullptr, 0, A_d, A_h,
+                                    Nbytes, hipMemcpyHostToDevice));
+  HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_B, graph, nullptr, 0, B_d, B_h,
+                                    Nbytes, hipMemcpyHostToDevice));
+  HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_B2Copies, graph, nullptr,
+                                    0, B_d, C_h,
+                                    Nbytes, hipMemcpyHostToDevice));
+
+  HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_C, graph, nullptr, 0, B_h, C_d,
+                                   Nbytes, hipMemcpyDeviceToHost));
+
+  void* kernelArgs2[] = {&A_d, &B_d, &C_d, reinterpret_cast<void *>(&NElem)};
+  kernelNodeParams.func = reinterpret_cast<void *>(HipTest::vectorADD<int>);
+  kernelNodeParams.gridDim = dim3(blocks);
+  kernelNodeParams.blockDim = dim3(threadsPerBlock);
+  kernelNodeParams.sharedMemBytes = 0;
+  kernelNodeParams.kernelParams = reinterpret_cast<void**>(kernelArgs2);
+  kernelNodeParams.extra = nullptr;
+  HIP_CHECK(hipGraphAddKernelNode(&kernel_vecAdd, graph, nullptr, 0,
+                                                        &kernelNodeParams));
+
+  // Create dependencies
+  HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_A, &kernel_vecAdd, 1));
+  HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_B, &kernel_vecAdd, 1));
+  HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_B2Copies, &kernel_vecAdd,
+                                    1));
+  HIP_CHECK(hipGraphAddDependencies(graph, &kernel_vecAdd, &memcpyD2H_C, 1));
+
+  // Destroy one of the dependency node
+  HIP_CHECK(hipGraphDestroyNode(memcpyH2D_B));
+
+  // Instantiate and launch the graph
+  HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
+  HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
+  HIP_CHECK(hipStreamSynchronize(streamForGraph));
+
+  // Verify graph execution result
+  HipTest::checkVectorADD(A_h, C_h, B_h, N);
+
+  HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
+  HIP_CHECK(hipGraphExecDestroy(graphExec));
+  HIP_CHECK(hipStreamDestroy(streamForGraph));
+  HIP_CHECK(hipGraphDestroy(graph));
+}

tests/catch/unit/graph/hipGraphExecHostNodeSetParams.cc

@@ -0,0 +1,276 @@
+/*
+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, INCLUDING 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 ANY 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.
+*/
+
+/**
+Testcase Scenarios of hipGraphExecHostNodeSetParams API:
+
+Functional:
+1. Creates graph, Adds HostNode, update hostNode params using hipGraphExecHostNodeSetParams API
+   and validates the result
+2. Create graph, Add Graphnodes and clones the graph. Add Hostnode to the cloned graph, update
+   hostNode params using hipGraphExecHostNodeSetParams API and validate the result
+
+Negative:
+
+1) Pass hGraphExec as nullptr and verify api doen't crash, returns error code.
+2) Pass node as nullptr and verify api doen't crash, returns error code.
+3) Pass pNodeParams as nullptr and verify api doesn't crash, returns error code.
+3) Pass hipHostNodeParams::hipHostFn_t as nullptr and verify api doesn't crash, returns error code.
+4) Pass unintialized host params and verify api doesn't crash, returns error code.
+5) Pass unintialized graph and verify api doesn't crash, returns error code.
+6) Pass nullptr to hostfunc and verify api doesn't crash, returns error code.
+*/
+
+#include <hip_test_common.hh>
+#include <hip_test_checkers.hh>
+
+#define SIZE 1024
+
+void callbackfunc(void *A_h) {
+  int *A = reinterpret_cast<int *>(A_h);
+  for (int i = 0; i < SIZE; i++) {
+     A[i] = i;
+  }
+}
+
+void callbackfunc_setparams(void *B_h) {
+  int *B = reinterpret_cast<int *>(B_h);
+  for (int i = 0; i < SIZE; i++) {
+     B[i] = i * i;
+  }
+}
+
+/*
+This testcase verifies the negative scenarios of
+hipGraphExecHostNodeSetParams API
+*/
+TEST_CASE("Unit_hipGraphExecHostNodeSetParams_Negative") {
+constexpr size_t N = 1024;
+  constexpr size_t Nbytes = N * sizeof(int);
+  hipGraph_t graph;
+  hipGraphExec_t graphExec;
+  int *A_d{nullptr}, *C_d{nullptr};
+  int *A_h{nullptr}, *C_h{nullptr};
+  HipTest::initArrays<int>(&A_d, nullptr, &C_d,
+                           &A_h, nullptr, &C_h, N, false);
+
+  HIP_CHECK(hipGraphCreate(&graph, 0));
+  hipGraphNode_t memcpyH2D_A, memcpyD2H_AC, memcpyH2D_C;
+  hipStream_t streamForGraph;
+  HIP_CHECK(hipStreamCreate(&streamForGraph));
+  HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr,
+                                    0, A_d, A_h,
+                                    Nbytes, hipMemcpyHostToDevice));
+  HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_C, graph, nullptr,
+                                    0, C_d, C_h,
+                                    Nbytes, hipMemcpyHostToDevice));
+  HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_AC, graph, nullptr,
+                                    0, A_h, C_d,
+                                    Nbytes, hipMemcpyDeviceToHost));
+  hipGraphNode_t hostNode;
+  hipHostNodeParams hostParams = {0, 0};
+  hostParams.fn = callbackfunc;
+  hostParams.userData = A_h;
+  HIP_CHECK(hipGraphAddHostNode(&hostNode, graph,
+                                 nullptr,
+                                 0, &hostParams));
+
+  HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_A,
+                                    &memcpyD2H_AC, 1));
+  HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_C,
+                                    &memcpyD2H_AC, 1));
+  HIP_CHECK(hipGraphAddDependencies(graph, &memcpyD2H_AC,
+                                    &hostNode, 1));
+
+  hipHostNodeParams sethostParams = {0, 0};
+  sethostParams.fn = callbackfunc_setparams;
+  sethostParams.userData = C_h;
+
+  // Instantiate and launch the graph
+  HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
+#if HT_NVIDIA
+  SECTION("Passing nullptr to graphExec") {
+    REQUIRE(hipGraphExecHostNodeSetParams(nullptr, hostNode, &sethostParams)
+                                      == hipErrorInvalidValue);
+  }
+
+  SECTION("Passing nullptr to hostParams") {
+    REQUIRE(hipGraphExecHostNodeSetParams(graphExec, hostNode, nullptr)
+                                      == hipErrorInvalidValue);
+  }
+#endif
+  SECTION("Passing nullptr to graph") {
+    REQUIRE(hipGraphExecHostNodeSetParams(graphExec, nullptr, &sethostParams)
+                                      == hipErrorInvalidValue);
+  }
+
+  SECTION("Passing nullptr to host func") {
+    sethostParams.fn = nullptr;
+    REQUIRE(hipGraphExecHostNodeSetParams(graphExec, hostNode, &sethostParams)
+                                      == hipErrorInvalidValue);
+  }
+
+
+  SECTION("Passing unintialized hostParams") {
+    hipHostNodeParams unintParams = {0, 0};
+    REQUIRE(hipGraphExecHostNodeSetParams(graphExec, hostNode, &unintParams)
+                                      == hipErrorInvalidValue);
+  }
+  HIP_CHECK(hipGraphDestroy(graph));
+}
+/*
+This testcase verifies hipGraphExecHostNodeSetParams API in cloned graph
+Creates graph, Add graph nodes and clone the graph
+Add HostNode to the cloned graph,update the host params using
+hipGraphExecHostNodeSetParams API and validates the result
+*/
+TEST_CASE("Unit_hipGraphExecHostNodeSetParams_ClonedGraphwithHostNode") {
+  constexpr size_t N = 1024;
+  constexpr size_t Nbytes = N * sizeof(int);
+  hipGraph_t graph;
+  hipGraphExec_t graphExec;
+  int *A_d{nullptr}, *C_d{nullptr};
+  int *A_h{nullptr}, *C_h{nullptr};
+  HipTest::initArrays<int>(&A_d, nullptr, &C_d,
+                           &A_h, nullptr, &C_h, N, false);
+
+  HIP_CHECK(hipGraphCreate(&graph, 0));
+  hipGraphNode_t memcpyH2D_A, memcpyH2D_C,
+                 memcpyD2H_AC;
+  hipStream_t streamForGraph;
+  HIP_CHECK(hipStreamCreate(&streamForGraph));
+  HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr,
+                                    0, A_d, A_h,
+                                    Nbytes, hipMemcpyHostToDevice));
+  HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_C, graph, nullptr,
+                                    0, C_d, C_h,
+                                    Nbytes, hipMemcpyHostToDevice));
+  HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_AC, graph, nullptr,
+                                    0, A_h, C_d,
+                                    Nbytes, hipMemcpyDeviceToHost));
+
+  hipGraph_t clonedgraph;
+  HIP_CHECK(hipGraphClone(&clonedgraph, graph));
+
+  hipGraphNode_t hostNode;
+  hipHostNodeParams hostParams = {0, 0};
+  hostParams.fn = callbackfunc;
+  hostParams.userData = A_h;
+  HIP_CHECK(hipGraphAddHostNode(&hostNode, clonedgraph,
+                                 nullptr,
+                                 0, &hostParams));
+
+  HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_A,
+                                    &memcpyD2H_AC, 1));
+  HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_C,
+                                    &memcpyD2H_AC, 1));
+
+  hipHostNodeParams sethostParams = {0, 0};
+  sethostParams.fn = callbackfunc_setparams;
+  sethostParams.userData = C_h;
+
+
+  // Instantiate and launch the cloned graph
+  HIP_CHECK(hipGraphInstantiate(&graphExec, clonedgraph, nullptr, nullptr, 0));
+  HIP_CHECK(hipGraphExecHostNodeSetParams(graphExec, hostNode, &sethostParams));
+  HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
+  HIP_CHECK(hipStreamSynchronize(streamForGraph));
+
+  // Verify execution result
+  for (size_t i = 0; i < N; i++) {
+    if (C_h[i] != static_cast<int>(i * i)) {
+      INFO("Validation failed i " <<  i  << "C_h[i] "<< C_h[i]);
+      REQUIRE(false);
+    }
+  }
+
+  HipTest::freeArrays<int>(A_d, nullptr, C_d, A_h, nullptr, C_h, false);
+  HIP_CHECK(hipGraphExecDestroy(graphExec));
+  HIP_CHECK(hipGraphDestroy(graph));
+  HIP_CHECK(hipStreamDestroy(streamForGraph));
+}
+
+/*
+This testcase verifies the following scenario
+Create graph, Adds host node to the graph,
+updates the host params using hipGraphExecHostNodeSetParams API
+and validates the result
+*/
+TEST_CASE("Unit_hipGraphExecHostNodeSetParams_BasicFunc") {
+  constexpr size_t N = 1024;
+  constexpr size_t Nbytes = N * sizeof(int);
+  hipGraph_t graph;
+  hipGraphExec_t graphExec;
+  int *A_d{nullptr}, *C_d{nullptr};
+  int *A_h{nullptr}, *C_h{nullptr};
+  HipTest::initArrays<int>(&A_d, nullptr, &C_d,
+                           &A_h, nullptr, &C_h, N, false);
+
+  HIP_CHECK(hipGraphCreate(&graph, 0));
+  hipGraphNode_t memcpyH2D_A, memcpyD2H_AC, memcpyH2D_C;
+  hipStream_t streamForGraph;
+  HIP_CHECK(hipStreamCreate(&streamForGraph));
+  HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr,
+                                    0, A_d, A_h,
+                                    Nbytes, hipMemcpyHostToDevice));
+  HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_C, graph, nullptr,
+                                    0, C_d, C_h,
+                                    Nbytes, hipMemcpyHostToDevice));
+  HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_AC, graph, nullptr,
+                                    0, A_h, C_d,
+                                    Nbytes, hipMemcpyDeviceToHost));
+  hipGraphNode_t hostNode;
+  hipHostNodeParams hostParams = {0, 0};
+  hostParams.fn = callbackfunc;
+  hostParams.userData = A_h;
+  HIP_CHECK(hipGraphAddHostNode(&hostNode, graph,
+                                 nullptr,
+                                 0, &hostParams));
+
+  HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_A,
+                                    &memcpyD2H_AC, 1));
+  HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_C,
+                                    &memcpyD2H_AC, 1));
+  HIP_CHECK(hipGraphAddDependencies(graph, &memcpyD2H_AC,
+                                    &hostNode, 1));
+
+  hipHostNodeParams sethostParams = {0, 0};
+  sethostParams.fn = callbackfunc_setparams;
+  sethostParams.userData = C_h;
+
+  // Instantiate and launch the graph
+  HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
+
+  HIP_CHECK(hipGraphExecHostNodeSetParams(graphExec, hostNode, &sethostParams));
+  HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
+  HIP_CHECK(hipStreamSynchronize(streamForGraph));
+
+  // Verify execution result
+  for (size_t i = 0; i < N; i++) {
+    if (C_h[i] != static_cast<int >(i * i)) {
+      INFO("Validation failed i " <<  i  << "C_h[i] "<< C_h[i]);
+      REQUIRE(false);
+    }
+  }
+
+  HipTest::freeArrays<int>(A_d, nullptr, C_d, A_h, nullptr, C_h, false);
+  HIP_CHECK(hipGraphExecDestroy(graphExec));
+  HIP_CHECK(hipGraphDestroy(graph));
+  HIP_CHECK(hipStreamDestroy(streamForGraph));
+}

tests/catch/unit/graph/hipGraphExecMemsetNodeSetParams.cc

@@ -0,0 +1,201 @@
+/*
+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, INCLUDING 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 ANY 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.
+*/
+
+/**
+Testcase Scenarios :
+Functional-
+1) Instantiate a graph with memset node, obtain executable graph and update the
+   hipMemsetParams node params with set. Make sure they are taking effect.
+Negative-
+1) Pass hGraphExec as nullptr and verify api returns error code.
+2) Pass graph node as nullptr and verify api returns error code.
+3) Pass different hipGraphNode_t which was not used in graphExec and verify api returns error code.
+4) Pass Pass different Graph which was not used in graphExec and verify api returns error code.
+5) Pass pNodeParams as nullptr and verify api returns error code.
+6) Pass pNodeParams as empty structure object and verify api returns error code.
+7) Pass hipMemsetParams::dst as nullptr, api should return error code.
+8) Pass hipMemsetParams::element size other than 1, 2, or 4 and check api should return error code.
+9) Pass hipMemsetParams::height as zero and check api should return error code.
+*/
+
+#include <hip_test_common.hh>
+#include <hip_test_checkers.hh>
+
+/* Test verifies hipGraphExecMemsetNodeSetParams API Negative scenarios.
+ */
+TEST_CASE("Unit_hipGraphExecMemsetNodeSetParams_Negative") {
+  constexpr size_t N = 1024;
+  constexpr size_t Nbytes = N * sizeof(char);
+  constexpr size_t val = 0;
+  char *devData, *hOutputData;
+
+  HIP_CHECK(hipMalloc(&devData, Nbytes));
+  hOutputData = reinterpret_cast<char *>(malloc(Nbytes));
+  REQUIRE(hOutputData != nullptr);
+  memset(hOutputData, 0,  Nbytes);
+
+  hipGraph_t graph;
+  hipError_t ret;
+  hipGraphExec_t graphExec;
+  hipStream_t streamForGraph;
+  hipGraphNode_t memsetNode;
+
+  HIP_CHECK(hipGraphCreate(&graph, 0));
+  HIP_CHECK(hipStreamCreate(&streamForGraph));
+
+  hipMemsetParams mParams{};
+  memset(&mParams, 0, sizeof(mParams));
+  mParams.dst = reinterpret_cast<void*>(devData);
+  mParams.value = val;
+  mParams.pitch = 0;
+  mParams.elementSize = sizeof(char);
+  mParams.width = Nbytes;
+  mParams.height = 1;
+  HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, &mParams));
+
+  std::vector<hipGraphNode_t> dependencies;
+  dependencies.push_back(memsetNode);
+
+  HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
+
+  SECTION("Pass hGraphExec as nullptr") {
+    ret = hipGraphExecMemsetNodeSetParams(nullptr, memsetNode, &mParams);
+    REQUIRE(hipErrorInvalidValue == ret);
+  }
+  SECTION("Pass hGraphNode as nullptr") {
+    ret = hipGraphExecMemsetNodeSetParams(graphExec, nullptr, &mParams);
+    REQUIRE(hipErrorInvalidValue == ret);
+  }
+  SECTION("Pass different hGraphNode which was not used in graphExec") {
+    hipGraphNode_t memsetNode1{};
+    ret = hipGraphExecMemsetNodeSetParams(graphExec, memsetNode1, &mParams);
+    REQUIRE(hipErrorInvalidValue == ret);
+  }
+  SECTION("Pass different Graph which was not used in graphExec") {
+    hipGraph_t graph1;
+    HIP_CHECK(hipGraphCreate(&graph1, 0));
+    HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph1, nullptr, 0, &mParams));
+    ret = hipGraphExecMemsetNodeSetParams(graphExec, memsetNode, &mParams);
+    REQUIRE(hipErrorInvalidValue == ret);
+    HIP_CHECK(hipGraphDestroy(graph1));
+  }
+  SECTION("Pass pNodeParams as nullptr") {
+    ret = hipGraphExecMemsetNodeSetParams(graphExec, memsetNode, nullptr);
+    REQUIRE(hipErrorInvalidValue == ret);
+  }
+#if HT_NVIDIA
+  SECTION("Pass pNodeParams as empty structure object") {
+    hipMemsetParams mParmTemp{};
+    ret = hipGraphExecMemsetNodeSetParams(graphExec, memsetNode, &mParmTemp);
+    REQUIRE(hipErrorInvalidValue == ret);
+  }
+#endif
+  SECTION("Pass hipMemsetParams::dst as nullptr") {
+    mParams.dst = nullptr;
+    ret = hipGraphExecMemsetNodeSetParams(graphExec, memsetNode, &mParams);
+    REQUIRE(hipErrorInvalidValue == ret);
+  }
+#if HT_NVIDIA
+  SECTION("Pass hipMemsetParams::element size other than 1, 2, or 4") {
+    mParams.dst = reinterpret_cast<void*>(devData);
+    mParams.elementSize = 9;
+    ret = hipGraphExecMemsetNodeSetParams(graphExec, memsetNode, &mParams);
+    REQUIRE(hipErrorInvalidValue == ret);
+  }
+  SECTION("Pass hipMemsetParams::height as zero") {
+    mParams.elementSize = sizeof(char);
+    mParams.height = 0;
+    ret = hipGraphExecMemsetNodeSetParams(graphExec, memsetNode, &mParams);
+    REQUIRE(hipErrorInvalidValue == ret);
+  }
+#endif
+
+  free(hOutputData);
+  HIP_CHECK(hipFree(devData));
+  HIP_CHECK(hipGraphExecDestroy(graphExec));
+  HIP_CHECK(hipGraphDestroy(graph));
+  HIP_CHECK(hipStreamDestroy(streamForGraph));
+}
+
+/* Test verifies hipGraphExecMemsetNodeSetParams API Functional scenarios.
+ */
+TEST_CASE("Unit_hipGraphExecMemsetNodeSetParams_Functional") {
+  constexpr size_t N = 1024;
+  constexpr size_t Nbytes = N * sizeof(char);
+  constexpr size_t val = 0;
+  constexpr size_t updateVal = 2;
+  char *devData, *devData1, *hOutputData, *hOutputData1;
+
+  HIP_CHECK(hipMalloc(&devData, Nbytes));
+  HIP_CHECK(hipMalloc(&devData1, Nbytes));
+  hOutputData = reinterpret_cast<char *>(malloc(Nbytes));
+  REQUIRE(hOutputData != nullptr);
+  memset(hOutputData, updateVal,  Nbytes);
+  hOutputData1 = reinterpret_cast<char *>(malloc(Nbytes));
+  REQUIRE(hOutputData1 != nullptr);
+  memset(hOutputData1, 0,  Nbytes);
+
+  hipGraph_t graph;
+  hipGraphExec_t graphExec;
+  hipStream_t streamForGraph;
+  hipGraphNode_t memsetNode;
+
+  HIP_CHECK(hipGraphCreate(&graph, 0));
+  HIP_CHECK(hipStreamCreate(&streamForGraph));
+
+  hipMemsetParams memsetParams{};
+  memset(&memsetParams, 0, sizeof(memsetParams));
+  memsetParams.dst = reinterpret_cast<void*>(devData);
+  memsetParams.value = val;
+  memsetParams.pitch = 0;
+  memsetParams.elementSize = sizeof(char);
+  memsetParams.width = Nbytes;
+  memsetParams.height = 1;
+  HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0,
+                                  &memsetParams));
+
+  std::vector<hipGraphNode_t> dependencies;
+  dependencies.push_back(memsetNode);
+
+  HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
+
+  memset(&memsetParams, 0, sizeof(memsetParams));
+  memsetParams.dst = reinterpret_cast<void*>(devData1);
+  memsetParams.value = updateVal;
+  memsetParams.pitch = 0;
+  memsetParams.elementSize = sizeof(char);
+  memsetParams.width = Nbytes;
+  memsetParams.height = 1;
+
+  REQUIRE(hipSuccess == hipGraphExecMemsetNodeSetParams(graphExec, memsetNode,
+                                                        &memsetParams));
+  HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
+  HIP_CHECK(hipStreamSynchronize(streamForGraph));
+
+  HIP_CHECK(hipMemcpy(hOutputData1, devData1, Nbytes, hipMemcpyDeviceToHost));
+  HipTest::checkArray(hOutputData, hOutputData1, Nbytes, 1);
+
+  free(hOutputData);
+  free(hOutputData1);
+  HIP_CHECK(hipFree(devData));
+  HIP_CHECK(hipFree(devData1));
+  HIP_CHECK(hipGraphExecDestroy(graphExec));
+  HIP_CHECK(hipGraphDestroy(graph));
+  HIP_CHECK(hipStreamDestroy(streamForGraph));
+}

tests/catch/unit/graph/hipGraphGetNodes.cc

@@ -0,0 +1,219 @@
+/*
+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, INCLUDING 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 ANY 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.
+*/
+
+/**
+Testcase Scenarios
+------------------
+Functional ::
+1) Add nodes to graph and get nodes. Verify the added nodes are present in returned list.
+2) Pass nodes as nullptr and verify numNodes returns actual number of nodes added to graph.
+3) If numNodes passed is greater than the actual number of nodes, the remaining entries in nodes
+will be set to NULL, and the number of nodes actually obtained will be returned in numNodes.
+
+Argument Validation ::
+1) Pass graph as nullptr and verify api returns error code.
+2) Pass numNodes as nullptr and other params as valid values. Expect api to return error code.
+3) When there are no nodes in graph, expect numNodes to be set to zero.
+4) Pass numNodes less than actual number of nodes. Expect api to populate requested number of node entries
+and does update numNodes.
+*/
+
+#include <hip_test_common.hh>
+#include <hip_test_checkers.hh>
+#include <hip_test_kernels.hh>
+
+/**
+ * Functional Test for hipGraphGetNodes API fetching node list
+ */
+TEST_CASE("Unit_hipGraphGetNodes_Functional") {
+  constexpr size_t N = 1024;
+  constexpr size_t Nbytes = N * sizeof(int);
+  constexpr auto blocksPerCU = 6;  // to hide latency
+  constexpr auto threadsPerBlock = 256;
+  constexpr auto addlEntries = 4;
+  hipGraph_t graph;
+  hipGraphNode_t memcpyNode, kernelNode;
+  hipKernelNodeParams kernelNodeParams{};
+  hipStream_t streamForGraph;
+  int *A_d, *B_d, *C_d;
+  int *A_h, *B_h, *C_h;
+  std::vector<hipGraphNode_t> dependencies, nodelist;
+  hipGraphExec_t graphExec;
+  size_t NElem{N};
+
+  HIP_CHECK(hipStreamCreate(&streamForGraph));
+  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(&memcpyNode, graph, NULL, 0, A_d, A_h,
+                                   Nbytes, hipMemcpyHostToDevice));
+  dependencies.push_back(memcpyNode);
+  nodelist.push_back(memcpyNode);
+  HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, NULL, 0, B_d, B_h,
+                                   Nbytes, hipMemcpyHostToDevice));
+  dependencies.push_back(memcpyNode);
+  nodelist.push_back(memcpyNode);
+
+  void* kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast<void *>(&NElem)};
+  kernelNodeParams.func = reinterpret_cast<void *>(HipTest::vectorADD<int>);
+  kernelNodeParams.gridDim = dim3(blocks);
+  kernelNodeParams.blockDim = dim3(threadsPerBlock);
+  kernelNodeParams.sharedMemBytes = 0;
+  kernelNodeParams.kernelParams = reinterpret_cast<void**>(kernelArgs);
+  kernelNodeParams.extra = nullptr;
+  HIP_CHECK(hipGraphAddKernelNode(&kernelNode, graph, dependencies.data(),
+                                  dependencies.size(), &kernelNodeParams));
+  dependencies.clear();
+  dependencies.push_back(kernelNode);
+  nodelist.push_back(kernelNode);
+  HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, dependencies.data(),
+                                    dependencies.size(), C_h, C_d,
+                                    Nbytes, hipMemcpyDeviceToHost));
+  nodelist.push_back(memcpyNode);
+
+  // Get numNodes by passing nodes as nullptr.
+  // verify : numNodes is set to actual number of nodes added
+  size_t numNodes{};
+  HIP_CHECK(hipGraphGetNodes(graph, nullptr, &numNodes));
+  INFO("Num of nodes returned by GetNodes : " << numNodes);
+  REQUIRE(numNodes == nodelist.size());
+
+  // Request for extra/additional nodes.
+  // verify : totNodes is reset to actual number of nodes
+  // verify : additional entries in nodes are set to nullptr
+  size_t totNodes = numNodes + addlEntries;
+  int numBytes = sizeof(hipGraphNode_t) * totNodes;
+  hipGraphNode_t* nodes = reinterpret_cast<hipGraphNode_t *>(malloc(numBytes));
+  REQUIRE(nodes != nullptr);
+  HIP_CHECK(hipGraphGetNodes(graph, nodes, &totNodes));
+  REQUIRE(totNodes == nodelist.size());
+  for (auto i = numNodes; i < numNodes + addlEntries; i++) {
+    REQUIRE(nodes[i] == nullptr);
+  }
+
+  // Verify added nodes are present in the node entries returned
+  for (auto Node : nodelist) {
+    bool found = false;
+    for (size_t i = 0; i < numNodes; i++) {
+      if (Node == nodes[i]) {
+        found = true;
+        break;
+      }
+    }
+
+    if (!found) {
+      INFO("Added node " << Node << " not present in returned list");
+      REQUIRE(false);
+    }
+  }
+
+  // Instantiate and launch the graph
+  HIP_CHECK(hipGraphInstantiate(&graphExec, graph, NULL, NULL, 0));
+  HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
+  HIP_CHECK(hipStreamSynchronize(streamForGraph));
+
+  // 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(streamForGraph));
+  free(nodes);
+}
+
+/**
+ * Test performs api parameter validation by passing various values
+ * as input and output parameters and validates the behavior.
+ * Test will include both negative and positive scenarios.
+ */
+TEST_CASE("Unit_hipGraphGetNodes_ParamValidation") {
+  hipStream_t stream{nullptr};
+  hipGraph_t graph{nullptr};
+  constexpr unsigned blocks = 512;
+  constexpr unsigned threadsPerBlock = 256;
+  constexpr size_t N = 1000000;
+  size_t Nbytes = N * sizeof(float), numNodes{};
+  float *A_d, *C_d;
+  float *A_h, *C_h;
+  A_h = reinterpret_cast<float*>(malloc(Nbytes));
+  C_h = reinterpret_cast<float*>(malloc(Nbytes));
+  REQUIRE(A_h != nullptr);
+  REQUIRE(C_h != nullptr);
+  HIP_CHECK(hipMalloc(&A_d, Nbytes));
+  HIP_CHECK(hipMalloc(&C_d, Nbytes));
+  REQUIRE(A_d != nullptr);
+  REQUIRE(C_d != nullptr);
+
+  HIP_CHECK(hipStreamCreate(&stream));
+  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));
+  HIP_CHECK(hipStreamEndCapture(stream, &graph));
+  HIP_CHECK(hipGraphGetNodes(graph, nullptr, &numNodes));
+  INFO("Num of nodes returned by GetNodes : " << numNodes);
+
+  int numBytes = sizeof(hipGraphNode_t) * numNodes;
+  hipGraphNode_t* nodes = reinterpret_cast<hipGraphNode_t *>(malloc(numBytes));
+  REQUIRE(nodes != nullptr);
+
+  SECTION("graph as nullptr") {
+    hipError_t ret = hipGraphGetNodes(nullptr, nodes, &numNodes);
+    REQUIRE(ret == hipErrorInvalidValue);
+  }
+
+  SECTION("numNodes as nullptr") {
+    hipError_t ret = hipGraphGetNodes(graph, nodes, nullptr);
+    REQUIRE(ret == hipErrorInvalidValue);
+  }
+
+  SECTION("no nodes in graph") {
+    hipGraph_t emptyGraph{};
+    HIP_CHECK(hipGraphCreate(&emptyGraph, 0));
+    HIP_CHECK(hipGraphGetNodes(emptyGraph, nullptr, &numNodes));
+    REQUIRE(numNodes == 0);
+  }
+
+  SECTION("numNodes less than actual number of nodes") {
+    size_t numPartNodes = numNodes - 1;
+    hipGraphNodeType nodeType;
+    HIP_CHECK(hipGraphGetNodes(graph, nodes, &numPartNodes));
+
+    // verify numPartNodes is unchanged
+    REQUIRE(numPartNodes == numNodes - 1);
+    // verify partial node list returned has valid nodes
+    for (size_t i = 0; i < numPartNodes; i++) {
+      HIP_CHECK(hipGraphNodeGetType(nodes[i], &nodeType));
+      REQUIRE(nodeType >= 0);
+      REQUIRE(nodeType < hipGraphNodeTypeCount);
+    }
+  }
+
+  HIP_CHECK(hipGraphDestroy(graph));
+  HIP_CHECK(hipStreamDestroy(stream));
+  free(A_h);
+  free(C_h);
+  free(nodes);
+  HIP_CHECK(hipFree(A_d));
+  HIP_CHECK(hipFree(C_d));
+}

tests/catch/unit/graph/hipGraphGetRootNodes.cc

@@ -0,0 +1,239 @@
+/*
+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, INCLUDING 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 ANY 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.
+*/
+
+/**
+Testcase Scenarios
+------------------
+Functional ::
+ 1) Add nodes to graph with and without dependencies, verify the api returns list of
+ root nodes (i.e., nodes without dependencies).
+ 2) Pass nodes as nullptr and verify api returns actual number of root nodes added to graph.
+ 3) If NumRootNodes passed is greater than the actual number of root nodes, the remaining entries in
+ nodes list will be set to NULL, and the number of nodes actually obtained will be returned in NumRootNodes.
+
+Argument Validation ::
+ 1) Pass graph as nullptr and verify api returns error code.
+ 2) Pass numRootNodes as nullptr and other params as valid values. Expect api to return error code.
+ 3) When there are no nodes in graph, expect numRootNodes to be set to zero.
+ 4) Pass numRootNodes less than actual number of nodes. Expect api to populate requested number of node entries
+ and does update numRootNodes.
+*/
+
+#include <hip_test_common.hh>
+#include <hip_test_checkers.hh>
+#include <hip_test_kernels.hh>
+
+/**
+ * Functional Test for API fetching root node list
+ */
+TEST_CASE("Unit_hipGraphGetRootNodes_Functional") {
+  constexpr size_t N = 1024;
+  constexpr size_t Nbytes = N * sizeof(int);
+  constexpr auto blocksPerCU = 6;  // to hide latency
+  constexpr auto threadsPerBlock = 256;
+  constexpr auto addlEntries = 5;
+  hipGraph_t graph;
+
+  hipGraphNode_t memcpyNode, kernelNode;
+  hipKernelNodeParams kernelNodeParams{};
+  hipStream_t streamForGraph;
+  int *A_d, *B_d, *C_d;
+  int *A_h, *B_h, *C_h;
+  std::vector<hipGraphNode_t> dependencies, rootnodelist;
+  hipGraphExec_t graphExec;
+  size_t NElem{N};
+
+  HIP_CHECK(hipStreamCreate(&streamForGraph));
+  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(&memcpyNode, graph, NULL, 0, A_d, A_h,
+                                   Nbytes, hipMemcpyHostToDevice));
+  dependencies.push_back(memcpyNode);
+  rootnodelist.push_back(memcpyNode);
+  HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, NULL, 0, B_d, B_h,
+                                   Nbytes, hipMemcpyHostToDevice));
+  dependencies.push_back(memcpyNode);
+  rootnodelist.push_back(memcpyNode);
+
+  void* kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast<void *>(&NElem)};
+  kernelNodeParams.func = reinterpret_cast<void *>(HipTest::vectorADD<int>);
+  kernelNodeParams.gridDim = dim3(blocks);
+  kernelNodeParams.blockDim = dim3(threadsPerBlock);
+  kernelNodeParams.sharedMemBytes = 0;
+  kernelNodeParams.kernelParams = reinterpret_cast<void**>(kernelArgs);
+  kernelNodeParams.extra = nullptr;
+  HIP_CHECK(hipGraphAddKernelNode(&kernelNode, graph, dependencies.data(),
+                                  dependencies.size(), &kernelNodeParams));
+  dependencies.clear();
+  dependencies.push_back(kernelNode);
+  HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, dependencies.data(),
+                                    dependencies.size(), C_h, C_d,
+                                    Nbytes, hipMemcpyDeviceToHost));
+
+  // Get numRootNodes by passing rootnodes list as nullptr.
+  // verify : numRootNodes is set to actual number of root nodes added
+  size_t numRootNodes{};
+  HIP_CHECK(hipGraphGetRootNodes(graph, nullptr, &numRootNodes));
+  INFO("Num of nodes returned by GetRootNodes : " << numRootNodes);
+  REQUIRE(numRootNodes == rootnodelist.size());
+
+  // Request for extra/additional nodes.
+  // verify : totNodes is reset to actual number of root nodes present
+  // verify : additional entries in rootnodes list are set to nullptr
+  size_t totNodes = numRootNodes + addlEntries;
+  int numBytes = sizeof(hipGraphNode_t) * totNodes;
+  hipGraphNode_t* rootnodes =
+                 reinterpret_cast<hipGraphNode_t *>(malloc(numBytes));
+  REQUIRE(rootnodes != nullptr);
+  HIP_CHECK(hipGraphGetRootNodes(graph, rootnodes, &totNodes));
+  REQUIRE(totNodes == rootnodelist.size());
+  for (auto i = numRootNodes; i < numRootNodes + addlEntries; i++) {
+    REQUIRE(rootnodes[i] == nullptr);
+  }
+
+  // Verify added nodes(without dependencies) are present
+  // in the root nodes fetched.
+  for (auto Node : rootnodelist) {
+    bool found = false;
+    for (size_t i = 0; i < numRootNodes; i++) {
+      if (Node == rootnodes[i]) {
+        found = true;
+        break;
+      }
+    }
+
+    if (!found) {
+      INFO("Returned root node " << Node << " not present in added list");
+      REQUIRE(false);
+    }
+  }
+
+  // Instantiate and launch the graph
+  HIP_CHECK(hipGraphInstantiate(&graphExec, graph, NULL, NULL, 0));
+  HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
+  HIP_CHECK(hipStreamSynchronize(streamForGraph));
+
+  // 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(streamForGraph));
+  free(rootnodes);
+}
+
+/**
+ * Test performs api parameter validation by passing various values
+ * as input and output parameters and validates the behavior.
+ * Test will include both negative and positive scenarios.
+ */
+TEST_CASE("Unit_hipGraphGetRootNodes_ParamValidation") {
+  hipStream_t stream1{nullptr}, stream2{nullptr}, mstream{nullptr};
+  hipEvent_t memsetEvent1, memsetEvent2, forkStreamEvent;
+  hipGraph_t graph{nullptr};
+  constexpr unsigned blocks = 512;
+  constexpr unsigned threadsPerBlock = 256;
+  constexpr size_t N = 1000000;
+  size_t Nbytes = N * sizeof(float), numRootNodes{};
+  float *A_d, *C_d;
+  float *A_h, *C_h;
+  A_h = reinterpret_cast<float*>(malloc(Nbytes));
+  C_h = reinterpret_cast<float*>(malloc(Nbytes));
+  REQUIRE(A_h != nullptr);
+  REQUIRE(C_h != nullptr);
+  HIP_CHECK(hipMalloc(&A_d, Nbytes));
+  HIP_CHECK(hipMalloc(&C_d, Nbytes));
+  REQUIRE(A_d != nullptr);
+  REQUIRE(C_d != nullptr);
+
+  HIP_CHECK(hipStreamCreate(&stream1));
+  HIP_CHECK(hipStreamCreate(&stream2));
+  HIP_CHECK(hipStreamCreate(&mstream));
+  HIP_CHECK(hipEventCreate(&memsetEvent1));
+  HIP_CHECK(hipEventCreate(&memsetEvent2));
+  HIP_CHECK(hipEventCreate(&forkStreamEvent));
+  HIP_CHECK(hipStreamBeginCapture(mstream, hipStreamCaptureModeGlobal));
+  HIP_CHECK(hipEventRecord(forkStreamEvent, mstream));
+  HIP_CHECK(hipStreamWaitEvent(stream1, forkStreamEvent, 0));
+  HIP_CHECK(hipStreamWaitEvent(stream2, forkStreamEvent, 0));
+  HIP_CHECK(hipMemsetAsync(A_d, 0, Nbytes, stream1));
+  HIP_CHECK(hipEventRecord(memsetEvent1, stream1));
+  HIP_CHECK(hipMemsetAsync(C_d, 0, Nbytes, stream2));
+  HIP_CHECK(hipEventRecord(memsetEvent2, stream2));
+  HIP_CHECK(hipStreamWaitEvent(mstream, memsetEvent1, 0));
+  HIP_CHECK(hipStreamWaitEvent(mstream, memsetEvent2, 0));
+  HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, mstream));
+  hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks),
+                              dim3(threadsPerBlock), 0, mstream, A_d, C_d, N);
+  HIP_CHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, mstream));
+  HIP_CHECK(hipStreamEndCapture(mstream, &graph));
+  HIP_CHECK(hipGraphGetRootNodes(graph, nullptr, &numRootNodes));
+  INFO("Num of nodes returned by GetRootNodes : " << numRootNodes);
+  int numBytes = sizeof(hipGraphNode_t) * numRootNodes;
+  hipGraphNode_t* nodes = reinterpret_cast<hipGraphNode_t *>(malloc(numBytes));
+  REQUIRE(nodes != nullptr);
+
+  SECTION("graph as nullptr") {
+    hipError_t ret = hipGraphGetRootNodes(nullptr, nodes, &numRootNodes);
+    REQUIRE(ret == hipErrorInvalidValue);
+  }
+
+  SECTION("numRootNodes as nullptr") {
+    hipError_t ret = hipGraphGetRootNodes(graph, nodes, nullptr);
+    REQUIRE(ret == hipErrorInvalidValue);
+  }
+
+  SECTION("no nodes in graph") {
+    hipGraph_t emptyGraph{};
+    HIP_CHECK(hipGraphCreate(&emptyGraph, 0));
+    HIP_CHECK(hipGraphGetRootNodes(emptyGraph, nullptr, &numRootNodes));
+    REQUIRE(numRootNodes == 0);
+  }
+
+  SECTION("numRootNodes less than actual number of nodes") {
+    size_t numPartNodes = numRootNodes - 1;
+    hipGraphNodeType nodeType;
+    HIP_CHECK(hipGraphGetRootNodes(graph, nodes, &numPartNodes));
+
+    // verify numPartNodes is unchanged
+    REQUIRE(numPartNodes == numRootNodes - 1);
+    // verify partial node list returned has valid nodes
+    for (size_t i = 0; i < numPartNodes; i++) {
+      HIP_CHECK(hipGraphNodeGetType(nodes[i], &nodeType));
+      REQUIRE(nodeType >= 0);
+      REQUIRE(nodeType < hipGraphNodeTypeCount);
+    }
+  }
+
+  HIP_CHECK(hipGraphDestroy(graph));
+  HIP_CHECK(hipStreamDestroy(mstream));
+  HIP_CHECK(hipStreamDestroy(stream1));
+  HIP_CHECK(hipStreamDestroy(stream2));
+  HIP_CHECK(hipEventDestroy(forkStreamEvent));
+  HIP_CHECK(hipEventDestroy(memsetEvent1));
+  HIP_CHECK(hipEventDestroy(memsetEvent2));
+  free(A_h);
+  free(C_h);
+  free(nodes);
+  HIP_CHECK(hipFree(A_d));
+  HIP_CHECK(hipFree(C_d));
+}

tests/catch/unit/graph/hipGraphInstantiateWithFlags.cc

@@ -0,0 +1,304 @@
+/*
+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, INCLUDING 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 ANY 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.
+*/
+
+/*
+hipGraphInstantiateWithFlags(hipGraphExec_t* pGraphExec, hipGraph_t graph, unsigned long long flags);
+Testcase Scenarios of hipGraphInstantiateWithFlags API:
+
+Negative:
+1) Pass nullptr to pGraphExec
+2) Pass nullptr to graph
+4) Pass invalid flag
+
+Functional:
+
+1) Create dependencies graph and instantiate the graph
+2) Create graph in one GPU device and instantiate, launch in peer GPU device
+3) Create stream capture graph and instantite the graph
+4) Create stream capture graph in one GPU device  and instantite the graph launch
+   in peer GPU device
+
+Mapping is missing for NVIDIA platform hence skipping the testcases
+*/
+
+
+#include <hip_test_common.hh>
+#include <hip_test_checkers.hh>
+#include <hip_test_kernels.hh>
+
+constexpr size_t N = 1000000;
+#if HT_AMD
+/* This test covers the negative scenarios of
+   hipGraphInstantiateWithFlags API */
+TEST_CASE("Unit_hipGraphInstantiateWithFlags_Negative") {
+#if HT_NVIDIA
+  SECTION("Passing nullptr pGraphExec") {
+    hipGraph_t graph;
+    HIP_CHECK(hipGraphCreate(&graph, 0));
+    REQUIRE(hipGraphInstantiateWithFlags(nullptr,
+                                         graph, 0) == hipErrorInvalidValue);
+  }
+
+  SECTION("Passing nullptr to graph") {
+    hipGraphExec_t graphExec;
+    REQUIRE(hipGraphInstantiateWithFlags(&graphExec,
+                                         nullptr, 0) == hipErrorInvalidValue);
+  }
+
+  SECTION("Passing Invalid flag") {
+    hipGraph_t graph;
+    HIP_CHECK(hipGraphCreate(&graph, 0));
+    hipGraphExec_t graphExec;
+    REQUIRE(hipGraphInstantiateWithFlags(&graphExec, graph, 10) != hipSuccess);
+  }
+#endif
+}
+/*
+This function verifies the following scenarios
+1. Creates dependency graph, Instantiates the graph with flags and verifies it
+2. Creates graph on one GPU-1 device and instantiates the graph on peer GPU device
+*/
+void GraphInstantiateWithFlags_DependencyGraph(bool ctxt_change = false) {
+  constexpr size_t N = 1024;
+  constexpr size_t Nbytes = N * sizeof(int);
+  constexpr auto blocksPerCU = 6;  // to hide latency
+  constexpr auto threadsPerBlock = 256;
+  hipGraph_t graph;
+  hipGraphNode_t memset_A, memset_B, memsetKer_C;
+  hipGraphNode_t memcpyH2D_A, memcpyH2D_B, memcpyD2H_C;
+  hipGraphNode_t kernel_vecAdd;
+  hipKernelNodeParams kernelNodeParams{};
+  int *A_d, *B_d, *C_d;
+  int *A_h, *B_h, *C_h;
+  hipGraphExec_t graphExec;
+  hipMemsetParams memsetParams{};
+  int memsetVal{};
+  size_t NElem{N};
+
+  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));
+
+  memset(&memsetParams, 0, sizeof(memsetParams));
+  memsetParams.dst = reinterpret_cast<void*>(A_d);
+  memsetParams.value = 0;
+  memsetParams.pitch = 0;
+  memsetParams.elementSize = sizeof(char);
+  memsetParams.width = Nbytes;
+  memsetParams.height = 1;
+  HIP_CHECK(hipGraphAddMemsetNode(&memset_A, graph, nullptr, 0,
+                                                              &memsetParams));
+
+  memset(&memsetParams, 0, sizeof(memsetParams));
+  memsetParams.dst = reinterpret_cast<void*>(B_d);
+  memsetParams.value = 0;
+  memsetParams.pitch = 0;
+  memsetParams.elementSize = sizeof(char);
+  memsetParams.width = Nbytes;
+  memsetParams.height = 1;
+  HIP_CHECK(hipGraphAddMemsetNode(&memset_B, graph, nullptr, 0,
+                                                              &memsetParams));
+
+  void* kernelArgs1[] = {&C_d, &memsetVal, reinterpret_cast<void *>(&NElem)};
+  kernelNodeParams.func =
+                       reinterpret_cast<void *>(HipTest::memsetReverse<int>);
+  kernelNodeParams.gridDim = dim3(blocks);
+  kernelNodeParams.blockDim = dim3(threadsPerBlock);
+  kernelNodeParams.sharedMemBytes = 0;
+  kernelNodeParams.kernelParams = reinterpret_cast<void**>(kernelArgs1);
+  kernelNodeParams.extra = nullptr;
+  HIP_CHECK(hipGraphAddKernelNode(&memsetKer_C, graph, nullptr, 0,
+                                                        &kernelNodeParams));
+
+  HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, 0, A_d, A_h,
+                                   Nbytes, hipMemcpyHostToDevice));
+  HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_B, graph, nullptr, 0, B_d, B_h,
+                                   Nbytes, hipMemcpyHostToDevice));
+
+  HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_C, graph, nullptr, 0, C_h, C_d,
+                                   Nbytes, hipMemcpyDeviceToHost));
+
+  void* kernelArgs2[] = {&A_d, &B_d, &C_d, reinterpret_cast<void *>(&NElem)};
+  kernelNodeParams.func = reinterpret_cast<void *>(HipTest::vectorADD<int>);
+  kernelNodeParams.gridDim = dim3(blocks);
+  kernelNodeParams.blockDim = dim3(threadsPerBlock);
+  kernelNodeParams.sharedMemBytes = 0;
+  kernelNodeParams.kernelParams = reinterpret_cast<void**>(kernelArgs2);
+  kernelNodeParams.extra = nullptr;
+  HIP_CHECK(hipGraphAddKernelNode(&kernel_vecAdd, graph, nullptr, 0,
+                                                        &kernelNodeParams));
+
+  // Create dependencies
+  HIP_CHECK(hipGraphAddDependencies(graph, &memset_A, &memcpyH2D_A, 1));
+  HIP_CHECK(hipGraphAddDependencies(graph, &memset_B, &memcpyH2D_B, 1));
+  HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_A, &kernel_vecAdd, 1));
+  HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_B, &kernel_vecAdd, 1));
+  HIP_CHECK(hipGraphAddDependencies(graph, &memsetKer_C, &kernel_vecAdd, 1));
+  HIP_CHECK(hipGraphAddDependencies(graph, &kernel_vecAdd, &memcpyD2H_C, 1));
+
+  if (ctxt_change) {
+    HIP_CHECK(hipSetDevice(1));
+    HIP_CHECK(hipDeviceEnablePeerAccess(0, 0));
+  }
+  // Instantiate and launch the cloned graph
+  HIP_CHECK(hipGraphInstantiateWithFlags(&graphExec, graph, 0));
+  HIP_CHECK(hipGraphLaunch(graphExec, 0));
+
+  // 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));
+}
+
+/*
+This function verifies the following scenarios
+1. Creates stream capture graph, Instantiates the graph with flags and verifies it
+2. Creates graph on one GPU-1 device and instantiates the graph on peer GPU device
+*/
+void GraphInstantiateWithFlags_StreamCapture(bool deviceContextChg = false) {
+  float *A_d, *C_d;
+  float *A_h, *C_h;
+  size_t Nbytes = N * sizeof(float);
+  hipStream_t stream;
+  hipGraph_t graph{nullptr};
+  hipGraphExec_t graphExec{nullptr};
+
+  A_h = reinterpret_cast<float*>(malloc(Nbytes));
+  C_h = reinterpret_cast<float*>(malloc(Nbytes));
+  REQUIRE(A_h != nullptr);
+  REQUIRE(C_h != nullptr);
+
+  // Fill with Phi + i
+  for (size_t i = 0; i < N; i++) {
+      A_h[i] = 1.618f + i;
+  }
+  HIP_CHECK(hipMalloc(&A_d, Nbytes));
+  HIP_CHECK(hipMalloc(&C_d, Nbytes));
+  REQUIRE(A_d != nullptr);
+  REQUIRE(C_d != nullptr);
+  HIP_CHECK(hipGraphCreate(&graph, 0));
+
+
+  HIP_CHECK(hipStreamCreate(&stream));
+  constexpr unsigned blocks = 512;
+  constexpr unsigned threadsPerBlock = 256;
+
+  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));
+
+  HIP_CHECK(hipStreamEndCapture(stream, &graph));
+
+  if (deviceContextChg) {
+    HIP_CHECK(hipSetDevice(1));
+    HIP_CHECK(hipDeviceEnablePeerAccess(0, 0));
+  }
+
+  // Validate end capture is successful
+  REQUIRE(graph != nullptr);
+  HIP_CHECK(hipGraphInstantiateWithFlags(&graphExec, graph, 0));
+  REQUIRE(graphExec != nullptr);
+
+  HIP_CHECK(hipGraphLaunch(graphExec, stream));
+
+  HIP_CHECK(hipStreamSynchronize(stream));
+
+  HIP_CHECK(hipGraphExecDestroy(graphExec));
+  HIP_CHECK(hipGraphDestroy(graph));
+
+  // Validate the computation
+  for (size_t i = 0; i < N; i++) {
+    if (C_h[i] != A_h[i] * A_h[i]) {
+      UNSCOPED_INFO("A and C not matching at " << i);
+      REQUIRE(false);
+    }
+  }
+
+  HIP_CHECK(hipStreamDestroy(stream));
+  free(A_h);
+  free(C_h);
+  HIP_CHECK(hipFree(A_d));
+  HIP_CHECK(hipFree(C_d));
+}
+/*
+This testcase verifies hipGraphInstantiateWithFlags API
+by creating dependency graph and instantiate, launching and verifying
+the result
+*/
+TEST_CASE("Unit_hipGraphInstantiateWithFlags_DependencyGraph") {
+  GraphInstantiateWithFlags_DependencyGraph();
+}
+/*
+This testcase verifies hipGraphInstantiateWithFlags API
+by creating dependency graph on GPU-0 and instantiate, launching and verifying
+the result on GPU-1
+*/
+#if HT_NVIDIA
+TEST_CASE("Unit_hipGraphInstantiateWithFlags_DependencyGraphDeviceCtxtChg") {
+  int numDevices = 0;
+  int canAccessPeer = 0;
+  HIP_CHECK(hipGetDeviceCount(&numDevices));
+  if (numDevices > 1) {
+    HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1));
+    if (canAccessPeer) {
+      GraphInstantiateWithFlags_DependencyGraph(true);
+    } else {
+      SUCCEED("Machine does not seem to have P2P");
+    }
+  } else {
+    SUCCEED("skipped the testcase as no of devices is less than 2");
+  }
+}
+#endif
+/*
+This testcase verifies hipGraphInstantiateWithFlags API
+by creating capture graph and instantiate, launching and verifying
+the result
+*/
+TEST_CASE("Unit_hipGraphInstantiateWithFlags_StreamCapture") {
+  GraphInstantiateWithFlags_StreamCapture();
+}
+
+/*
+This testcase verifies hipGraphInstantiateWithFlags API
+by creating capture graph on GPU-0 and instantiate, launching and verifying
+the result on GPU-1
+*/
+TEST_CASE("Unit_hipGraphInstantiateWithFlags_StreamCaptureDeviceContextChg") {
+  int numDevices = 0;
+  int canAccessPeer = 0;
+  HIP_CHECK(hipGetDeviceCount(&numDevices));
+  if (numDevices > 1) {
+    HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1));
+    if (canAccessPeer) {
+      GraphInstantiateWithFlags_StreamCapture(true);
+    } else {
+      SUCCEED("Machine does not seem to have P2P");
+    }
+  } else {
+    SUCCEED("skipped the testcase as no of devices is less than 2");
+  }
+}
+#endif

tests/catch/unit/graph/hipGraphMemcpyNodeSetParamsToSymbol.cc

@@ -0,0 +1,107 @@
+#include <hip_test_common.hh>
+#include <hip_test_checkers.hh>
+#include <limits>
+#define SIZE 256
+
+__device__ int globalIn[SIZE], globalOut[SIZE];
+__device__ __constant__ int globalConst[SIZE];
+
+
+/* This testcase verifies negative scenarios of
+   hipGraphMemcpyNodeSetParamsToSymbol API */
+TEST_CASE("Unit_hipGraphMemcpyNodeSetParamsToSymbol_Negative") {
+  constexpr size_t Nbytes = SIZE * sizeof(int);
+  int *A_d{nullptr}, *B_d{nullptr};
+  int *A_h{nullptr}, *B_h{nullptr};
+  HipTest::initArrays<int>(&A_d, &B_d, nullptr,
+                           &A_h, &B_h, nullptr, SIZE, false);
+
+  hipGraph_t graph;
+  hipError_t ret;
+  hipGraphNode_t memcpyToSymbolNode, memcpyH2D_A;
+  std::vector<hipGraphNode_t> dependencies;
+  HIP_CHECK(hipGraphCreate(&graph, 0));
+
+  // Adding MemcpyNode
+  HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, 0, A_d, A_h,
+                                    Nbytes, hipMemcpyHostToDevice));
+  dependencies.push_back(memcpyH2D_A);
+
+  HIP_CHECK(hipGraphAddMemcpyNodeToSymbol(&memcpyToSymbolNode, graph,
+                                            dependencies.data(),
+                                            dependencies.size(),
+                                            HIP_SYMBOL(globalIn),
+                                            A_d, Nbytes, 0,
+                                            hipMemcpyDeviceToDevice));
+
+  SECTION("Pass GraphNode as nullptr") {
+    ret = hipGraphMemcpyNodeSetParamsToSymbol(nullptr,
+                                              HIP_SYMBOL(globalIn),
+                                              B_d, Nbytes, 0,
+                                              hipMemcpyDeviceToDevice);
+    REQUIRE(hipErrorInvalidValue == ret);
+  }
+  SECTION("Pass symbol ptr as nullptr") {
+    ret = hipGraphMemcpyNodeSetParamsToSymbol(memcpyToSymbolNode,
+                                              nullptr,
+                                              B_d, Nbytes, 0,
+                                              hipMemcpyDeviceToDevice);
+    REQUIRE(hipErrorInvalidSymbol == ret);
+  }
+  SECTION("Pass src ptr as nullptr") {
+    ret = hipGraphMemcpyNodeSetParamsToSymbol(memcpyToSymbolNode,
+                                              HIP_SYMBOL(globalIn),
+                                              nullptr, Nbytes, 0,
+                                              hipMemcpyDeviceToDevice);
+    REQUIRE(hipErrorInvalidValue == ret);
+  }
+  SECTION("Pass count as zero") {
+    ret = hipGraphMemcpyNodeSetParamsToSymbol(memcpyToSymbolNode,
+                                              HIP_SYMBOL(globalIn),
+                                              B_d, 0, 0,
+                                              hipMemcpyDeviceToDevice);
+    REQUIRE(hipErrorInvalidValue == ret);
+  }
+
+  SECTION("Pass count more than allocated size for source and dstn ptr") {
+    ret = hipGraphMemcpyNodeSetParamsToSymbol(memcpyToSymbolNode,
+                                              HIP_SYMBOL(globalIn),
+                                              B_d, Nbytes+8, 0,
+                                              hipMemcpyDeviceToDevice);
+    REQUIRE(hipErrorInvalidValue == ret);
+  }
+  SECTION("Pass offset+count greater than allocated size") {
+    ret = hipGraphMemcpyNodeSetParamsToSymbol(memcpyToSymbolNode,
+                                              HIP_SYMBOL(globalIn),
+                                              B_d, Nbytes, 10,
+                                              hipMemcpyDeviceToDevice);
+    REQUIRE(hipErrorInvalidValue == ret);
+  }
+  SECTION("Pass same symbol pointer as source ptr and destination ptr") {
+    ret = hipGraphMemcpyNodeSetParamsToSymbol(memcpyToSymbolNode,
+                                              HIP_SYMBOL(globalIn),
+                                              HIP_SYMBOL(globalIn),
+                                              Nbytes, 0,
+                                              hipMemcpyDeviceToDevice);
+    REQUIRE(hipErrorInvalidValue == ret);
+  }
+  SECTION("Pass 2 different symbol pointer as source ptr and dstn ptr") {
+    ret = hipGraphMemcpyNodeSetParamsToSymbol(memcpyToSymbolNode,
+                                              HIP_SYMBOL(globalIn),
+                                              HIP_SYMBOL(globalOut),
+                                              Nbytes, 0,
+                                              hipMemcpyDeviceToDevice);
+    REQUIRE(hipErrorInvalidValue == ret);
+  }
+  SECTION("Copy from host ptr to device ptr but pass kind as different") {
+    ret = hipGraphMemcpyNodeSetParamsToSymbol(memcpyToSymbolNode,
+                                              HIP_SYMBOL(globalIn),
+                                              A_h,
+                                              Nbytes, 0,
+                                              hipMemcpyDeviceToDevice);
+    REQUIRE(hipErrorInvalidValue == ret);
+  }
+
+  HipTest::freeArrays<int>(A_d, B_d, nullptr, A_h, B_h, nullptr, false);
+  HIP_CHECK(hipGraphDestroy(graph));
+}

tests/catch/unit/graph/hipGraphNodeFindInClone.cc

@@ -0,0 +1,241 @@
+/*
+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, INCLUDING 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 ANY 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.
+*/
+
+/*
+Testcase Scenarios of hipGraphNodeFindInClone API:
+
+Negative:
+
+1) Pass nullptr to graph node
+2) pass nullptr to original graph node
+3) pass nullptr to clonedGraph
+4) Pass original graph in place of the cloned graph
+5) Pass invalid originalNode
+6) Destroy the graph node in the original graph
+   and try to get the deleted graph node
+   from the cloned graph
+7) Clone the graph,Add node to Original graph
+   and try to find the original node in the cloned graph
+
+
+Functional:
+
+1) Get the graph node from the cloned graph corresponding to the original node
+2) Create and clone the graph, modify the original graph and clone the graph again,
+   then try to find the newly added graph node  from the cloned graph
+
+*/
+
+#include<hip/hip_runtime_api.h>
+#include <hip_test_common.hh>
+#include <hip_test_checkers.hh>
+#include <hip_test_kernels.hh>
+
+
+/* This test covers the negative scenarios of
+   hipGraphNodeFindInClone API */
+
+TEST_CASE("Unit_hipGraphNodeFindInClone_Negative") {
+  hipGraph_t graph;
+  hipGraph_t clonedgraph;
+  hipGraphNode_t graphnode, newnode;
+  hipGraphNode_t clonedgraphnode;
+  HIP_CHECK(hipGraphCreate(&graph, 0));
+
+
+  int *A_d, *A_h, *B_d, *B_h;
+  HipTest::initArrays<int>(&A_d, &B_d, nullptr, &A_h,
+                           &B_h, nullptr, 1024, false);
+  HIP_CHECK(hipGraphAddMemcpyNode1D(&graphnode, graph, nullptr, 0, A_d, A_h,
+                                    1024, hipMemcpyHostToDevice));
+  // Cloned the graph
+  HIP_CHECK(hipGraphClone(&clonedgraph, graph));
+  HIP_CHECK(hipGraphAddMemcpyNode1D(&newnode, graph, nullptr, 0, B_d, B_h,
+                                    1024, hipMemcpyHostToDevice));
+
+  SECTION("Passing nullptr to Cloned graph") {
+    REQUIRE(hipGraphNodeFindInClone(&clonedgraphnode, graphnode, nullptr)
+                                    == hipErrorInvalidValue);
+  }
+
+  SECTION("Passing nullptr to original graph") {
+    REQUIRE(hipGraphNodeFindInClone(nullptr, graphnode, clonedgraph)
+                                    == hipErrorInvalidValue);
+  }
+
+  SECTION("Passing nullptr to graph node") {
+    REQUIRE(hipGraphNodeFindInClone(&clonedgraphnode, nullptr, clonedgraph)
+                                    == hipErrorInvalidValue);
+  }
+#if HT_NVIDIA
+  SECTION("Pass uncloned graph") {
+    REQUIRE(hipGraphNodeFindInClone(&clonedgraphnode, graphnode, graph)
+                                    == hipErrorInvalidValue);
+  }
+
+  SECTION("Destroy the graph node and find in cloned graph") {
+    HIP_CHECK(hipGraphDestroyNode(graphnode));
+    REQUIRE(hipGraphNodeFindInClone(&clonedgraphnode, graphnode,
+                                    clonedgraph)
+        == hipErrorInvalidValue);
+  }
+#endif
+
+  SECTION("Pass invalid original graphnode") {
+    hipGraphNode_t unintialized_graphnode{nullptr};
+    REQUIRE(hipGraphNodeFindInClone(&clonedgraphnode, unintialized_graphnode,
+                                    graph)
+                                    == hipErrorInvalidValue);
+  }
+
+  SECTION("Find node in cloned graph which is only present in original graph") {
+    REQUIRE(hipGraphNodeFindInClone(&clonedgraphnode, newnode,
+                                    clonedgraph) == hipErrorInvalidValue);
+  }
+
+
+  HipTest::freeArrays<int>(A_d, B_d, nullptr,
+                           A_h, B_h, nullptr, false);
+  HIP_CHECK(hipGraphDestroy(graph));
+  HIP_CHECK(hipGraphDestroy(clonedgraph));
+}
+
+
+void hipGraphNodeFindInClone_Func(bool ModifyOrigGraph = false) {
+  constexpr size_t N = 1024;
+  constexpr size_t Nbytes = N * sizeof(int);
+  constexpr auto blocksPerCU = 6;  // to hide latency
+  constexpr auto threadsPerBlock = 256;
+  hipGraph_t graph, clonedgraph;
+  hipGraphNode_t memset_A, memset_B, memsetKer_C;
+  hipGraphNode_t memcpyH2D_A, memcpyH2D_B, memcpyD2H_C, memcpyD2D_C,
+                 memcpyD2H_C_new;
+  hipGraphNode_t kernel_vecAdd;
+  hipKernelNodeParams kernelNodeParams{};
+  int *A_d, *B_d, *C_d;
+  int *A_h, *B_h, *C_h;
+  hipMemsetParams memsetParams{};
+  int memsetVal{};
+  size_t NElem{N};
+
+  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));
+
+  memset(&memsetParams, 0, sizeof(memsetParams));
+  memsetParams.dst = reinterpret_cast<void*>(A_d);
+  memsetParams.value = 0;
+  memsetParams.pitch = 0;
+  memsetParams.elementSize = sizeof(char);
+  memsetParams.width = Nbytes;
+  memsetParams.height = 1;
+  HIP_CHECK(hipGraphAddMemsetNode(&memset_A, graph, nullptr, 0,
+                                                              &memsetParams));
+
+  memset(&memsetParams, 0, sizeof(memsetParams));
+  memsetParams.dst = reinterpret_cast<void*>(B_d);
+  memsetParams.value = 0;
+  memsetParams.pitch = 0;
+  memsetParams.elementSize = sizeof(char);
+  memsetParams.width = Nbytes;
+  memsetParams.height = 1;
+  HIP_CHECK(hipGraphAddMemsetNode(&memset_B, graph, nullptr, 0,
+                                                              &memsetParams));
+
+  void* kernelArgs1[] = {&C_d, &memsetVal, reinterpret_cast<void *>(&NElem)};
+  kernelNodeParams.func =
+                       reinterpret_cast<void *>(HipTest::memsetReverse<int>);
+  kernelNodeParams.gridDim = dim3(blocks);
+  kernelNodeParams.blockDim = dim3(threadsPerBlock);
+  kernelNodeParams.sharedMemBytes = 0;
+  kernelNodeParams.kernelParams = reinterpret_cast<void**>(kernelArgs1);
+  kernelNodeParams.extra = nullptr;
+  HIP_CHECK(hipGraphAddKernelNode(&memsetKer_C, graph, nullptr, 0,
+                                                        &kernelNodeParams));
+
+  HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, 0, A_d, A_h,
+                                   Nbytes, hipMemcpyHostToDevice));
+
+  HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_B, graph, nullptr, 0, B_d, B_h,
+                                   Nbytes, hipMemcpyHostToDevice));
+
+  HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_C, graph, nullptr, 0, C_h, C_d,
+                                   Nbytes, hipMemcpyDeviceToHost));
+
+  void* kernelArgs2[] = {&A_d, &B_d, &C_d, reinterpret_cast<void *>(&NElem)};
+  kernelNodeParams.func = reinterpret_cast<void *>(HipTest::vectorADD<int>);
+  kernelNodeParams.gridDim = dim3(blocks);
+  kernelNodeParams.blockDim = dim3(threadsPerBlock);
+  kernelNodeParams.sharedMemBytes = 0;
+  kernelNodeParams.kernelParams = reinterpret_cast<void**>(kernelArgs2);
+  kernelNodeParams.extra = nullptr;
+  HIP_CHECK(hipGraphAddKernelNode(&kernel_vecAdd, graph, nullptr, 0,
+                                                        &kernelNodeParams));
+
+  // Create dependencies
+  HIP_CHECK(hipGraphAddDependencies(graph, &memset_A, &memcpyH2D_A, 1));
+  HIP_CHECK(hipGraphAddDependencies(graph, &memset_B, &memcpyH2D_B, 1));
+  HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_A, &kernel_vecAdd, 1));
+  HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_B, &kernel_vecAdd, 1));
+  HIP_CHECK(hipGraphAddDependencies(graph, &memsetKer_C, &kernel_vecAdd, 1));
+  HIP_CHECK(hipGraphAddDependencies(graph, &kernel_vecAdd, &memcpyD2H_C, 1));
+
+
+  if (ModifyOrigGraph) {
+    // Cloned the graph
+    HIP_CHECK(hipGraphClone(&clonedgraph, graph));
+    // Modify Original graph by adding new dependency
+    HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2D_C, graph, nullptr, 0,
+                                      C_d, B_d,
+                                      Nbytes, hipMemcpyDeviceToHost));
+
+    HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_C_new, graph, nullptr, 0,
+                                      C_h, C_d,
+                                      Nbytes, hipMemcpyDeviceToHost));
+
+    HIP_CHECK(hipGraphAddDependencies(graph, &kernel_vecAdd, &memcpyD2D_C, 1));
+    HIP_CHECK(hipGraphAddDependencies(graph, &memcpyD2D_C,
+                                      &memcpyD2H_C_new, 1));
+  }
+  // Cloned the graph
+  HIP_CHECK(hipGraphClone(&clonedgraph, graph));
+  hipGraphNode_t clonedgraphnode;
+  if (ModifyOrigGraph) {
+    REQUIRE(hipGraphNodeFindInClone(&clonedgraphnode,
+                                    memcpyD2H_C_new, clonedgraph)
+                                    == hipSuccess);
+  } else {
+    REQUIRE(hipGraphNodeFindInClone(&clonedgraphnode,
+                                    memcpyH2D_A, clonedgraph)
+                                    == hipSuccess);
+  }
+  HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
+  HIP_CHECK(hipGraphDestroy(graph));
+  HIP_CHECK(hipGraphDestroy(clonedgraph));
+}
+
+TEST_CASE("Unit_hipGraphNodeFindInClone_Functional") {
+  SECTION("hipGraphNodeFindInClone Basic Functionality") {
+    hipGraphNodeFindInClone_Func();
+  }
+  SECTION("hipGraphNodeFindInClone Modify Original graph") {
+    hipGraphNodeFindInClone_Func(true);
+  }
+}

tests/catch/unit/memory/CMakeLists.txt

@@ -35,6 +35,9 @@ set(TEST_SRC
     hipMemPrefetchAsyncExtTsts.cc
     hipMemAdviseMmap.cc
     hipMemCoherencyTst.cc
+    hipMallocManaged.cc
+    hipMemRangeGetAttribute.cc
+    hipHmmOvrSubscriptionTst.cc
 )
 else()
 set(TEST_SRC
@@ -69,6 +72,9 @@ set(TEST_SRC
     hipMallocManagedFlagsTst.cc
     hipMemPrefetchAsyncExtTsts.cc
     hipMemAdviseMmap.cc
+    hipMallocManaged.cc
+    hipMemRangeGetAttribute.cc
+    hipHmmOvrSubscriptionTst.cc
 )
 endif()
 

tests/catch/unit/memory/hipHmmOvrSubscriptionTst.cc

@@ -0,0 +1,213 @@
+/*
+Copyright (c) 2021-Present 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.
+*/
+
+/* Test Case Description: This test case tests the working of OverSubscription
+   feature which is part of HMM.*/
+
+#include <hip_test_common.hh>
+#ifdef __linux__
+#include <sys/types.h>
+#include <sys/ipc.h>
+#include <sys/shm.h>
+#include <sys/stat.h>
+#include <stdio.h>
+#include <stdlib.h>
+#include <string.h>
+#include <unistd.h>
+#include <sys/wait.h>
+#endif
+#include <list>
+
+#define INIT_VAL 2.5
+#define NUM_ELMS 268435456  // 268435456 * 4 = 1GB
+#define ITERATIONS 10
+#define ONE_GB 1024 * 1024 * 1024
+
+static void GetTotGpuMem(int *TotMem);
+static void DisplayHmmFlgs(int *Signal);
+// Kernel function
+__global__ void Square(int n, float *x) {
+  int index = blockIdx.x * blockDim.x + threadIdx.x;
+  int stride = blockDim.x * gridDim.x;
+  for (int i = index; i < n; i += stride) {
+    x[i] = x[i] + 10;
+  }
+}
+
+static void OneGBMemTest(int dev) {
+  int DataMismatch = 0;
+  float *HmmAG = nullptr;
+  hipStream_t strm;
+  HIP_CHECK(hipStreamCreate(&strm));
+  // Testing hipMemAttachGlobal Flag
+  HIP_CHECK(hipMallocManaged(&HmmAG, NUM_ELMS * sizeof(float),
+                            hipMemAttachGlobal));
+
+  // Initializing HmmAG memory
+  for (int i = 0; i < NUM_ELMS; i++) {
+    HmmAG[i] = INIT_VAL;
+  }
+
+  int blockSize = 256;
+  int numBlocks = (NUM_ELMS + blockSize - 1) / blockSize;
+  dim3 dimGrid(numBlocks, 1, 1);
+  dim3 dimBlock(blockSize, 1, 1);
+  HIP_CHECK(hipSetDevice(dev));
+  for (int i = 0; i < ITERATIONS; ++i) {
+    Square<<<dimGrid, dimBlock, 0, strm>>>(NUM_ELMS, HmmAG);
+  }
+  HIP_CHECK(hipStreamSynchronize(strm));
+  for (int j = 0; j < NUM_ELMS; ++j) {
+    if (HmmAG[j] != (INIT_VAL + ITERATIONS * 10)) {
+      DataMismatch++;
+      break;
+    }
+  }
+  if (DataMismatch != 0) {
+    WARN("Data Mismatch observed when kernel launched on device: " << dev);
+    REQUIRE(false);
+  }
+  HIP_CHECK(hipFree(HmmAG));
+  HIP_CHECK(hipStreamDestroy(strm));
+}
+
+static void GetTotGpuMem(int *TotMem) {
+  size_t FreeMem, TotGpuMem;
+  HIP_CHECK(hipMemGetInfo(&FreeMem, &TotGpuMem));
+  TotMem[0] = (TotGpuMem/(ONE_GB));
+  TotMem[1] = 1;
+}
+
+static void DisplayHmmFlgs(int *Signal) {
+  int managed = 0;
+  WARN("The following are the attribute values related to HMM for"
+         " device 0:\n");
+  HIP_CHECK(hipDeviceGetAttribute(&managed,
+                    hipDeviceAttributeDirectManagedMemAccessFromHost, 0));
+  WARN("hipDeviceAttributeDirectManagedMemAccessFromHost: " << managed);
+  HIP_CHECK(hipDeviceGetAttribute(&managed,
+                                 hipDeviceAttributeConcurrentManagedAccess, 0));
+  WARN("hipDeviceAttributeConcurrentManagedAccess: " << managed);
+  HIP_CHECK(hipDeviceGetAttribute(&managed,
+                                 hipDeviceAttributePageableMemoryAccess, 0));
+  WARN("hipDeviceAttributePageableMemoryAccess: " << managed);
+  HIP_CHECK(hipDeviceGetAttribute(&managed,
+              hipDeviceAttributePageableMemoryAccessUsesHostPageTables, 0));
+  WARN("hipDeviceAttributePageableMemoryAccessUsesHostPageTables:"
+          << managed);
+
+  HIP_CHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeManagedMemory,
+                                  0));
+  WARN("hipDeviceAttributeManagedMemory: " << managed);
+
+  // Checking for Vega20 or MI100
+  hipDeviceProp_t prop;
+  HIP_CHECK(hipGetDeviceProperties(&prop, 0));
+  char *p = NULL;
+  p = strstr(prop.gcnArchName, "gfx906");
+  if (p) {
+    WARN("This system has MI60 gpu hence OverSubscription test will be");
+    WARN(" skipped");
+    Signal[2] = 1;
+  }
+  p = strstr(prop.gcnArchName, "gfx908");
+  if (p) {
+    WARN("This system has MI100 gpu hence OverSubscription test will be");
+    WARN(" skipped");
+    Signal[2] = 1;
+  }
+  Signal[1] = managed;
+  Signal[0] = 1;
+}
+
+TEST_CASE("Unit_HMM_OverSubscriptionTst") {
+  int HmmEnabled = 0;
+  // The following Shared Mem is to get Max GPU Mem
+  // The size requested is for three ints
+  // 1) To get Max GPU Mem in GB
+  // 2) To Signal parent that req. info is available to consume
+  // 3) To know if MI60 or MI100 gpu are there in the system
+  key_t key = ftok("shmTotMem", 66);
+  int shmid = shmget(key, (3 * sizeof(int)), 0666|IPC_CREAT);
+  int *TotGpuMem = reinterpret_cast<int*>(shmat(shmid, NULL, 0));
+  TotGpuMem[0] = 0; TotGpuMem[1] = 0;
+  // The following function DisplayHmmFlgs() displays the flag values related
+  // to HMM and also sends us ManagedMemory attribute value
+  if (fork() == 0) {
+    DisplayHmmFlgs(TotGpuMem);
+    exit(1);
+  }
+  while (TotGpuMem[0] == 0) {
+      sleep(2);
+  }
+  // The following if block will skip test if either of MI60 or MI100 is found
+  if (TotGpuMem[2] == 1) {
+    SUCCEED("Test is skipped!!");
+    REQUIRE(true);
+  } else {
+    HmmEnabled = TotGpuMem[1];
+
+    // Re-setting the shared memory values for further usage
+    TotGpuMem[0] = 0;
+    TotGpuMem[1] = 0;
+
+    std::list<pid_t> PidLst;
+    // The following function gets the MaxGpu memory in GBs and also launches
+    // OverSubscription test
+    if (HmmEnabled) {
+      if ((setenv("HSA_XNACK", "1", 1)) != 0) {
+        WARN("Unable to turn on HSA_XNACK, hence terminating the Test case!");
+        REQUIRE(false);
+      }
+      if (fork() == 0) {
+        GetTotGpuMem(TotGpuMem);
+      }
+      while (TotGpuMem[1] == 0) {
+        sleep(2);
+      }
+      int NumGB = TotGpuMem[0], TotalThreads = (NumGB + 10);
+      WARN("Launching " << TotalThreads);
+      WARN(" processes to test OverSubscription.");
+      pid_t pid;
+      for (int k = 0; k < TotalThreads; ++k) {
+        pid = fork();
+        PidLst.push_back(pid);
+        if (pid == 0) {
+          OneGBMemTest(0);
+          exit(10);
+        }
+      }
+    } else {
+      SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
+              "attribute. Hence skipping the testing with Pass result.\n");
+    }
+    int status;
+    for (pid_t pd : PidLst) {
+      waitpid(pd, &status, 0);
+      if (!(WIFEXITED(status))) {
+        REQUIRE(false);
+      }
+    }
+  }
+  shmdt(TotGpuMem);
+  shmctl(shmid, IPC_RMID, NULL);
+}

tests/catch/unit/memory/hipMallocManaged.cc

@@ -17,12 +17,11 @@
    THE SOFTWARE.
  */
 
-/*
-  List of Test cases:
-  1)  Unit_hipMallocManaged_Basic
-  2) Unit_hipMallocManaged_MultiSize
-  3) Unit_hipMallocManaged_MultiKrnlHmmAccess
-  4) Unit_hipMallocManaged_KrnlWth2MemTypes
+/* Test Case Description:
+   1) This testcase verifies the hipMallocManaged basic scenario - supported on
+     all devices
+   2) This testcase verifies the hipMallocManaged basic scenario - supported
+       only on HMM enabled devices
 */
 
 #include <hip_test_common.hh>
@@ -32,12 +31,6 @@
 
 
 // Kernel functions
-__global__ void KrnlWth2MemTypes(int *Hmm, int *Dptr, size_t n) {
-  size_t index = blockIdx.x * blockDim.x + threadIdx.x;
-  for (size_t i = index; i < n; i++) {
-    Hmm[i] = Dptr[i] + 10;
-  }
-}
 
 __global__ void KernelMul_MngdMem(int *Hmm, int *Dptr, size_t n) {
   size_t index = blockIdx.x * blockDim.x + threadIdx.x;
@@ -64,9 +57,6 @@ __global__ void KrnlWth2MemTypesC(unsigned char *Hmm, unsigned char *Dptr,
   }
 }
 
-// The following variable will be used to get the result of computation
-// from multiple threads
-static bool IfTestPassed = true;
 
 static int HmmAttrPrint() {
   int managed = 0;
@@ -93,62 +83,6 @@ static int HmmAttrPrint() {
 }
 
 
-static void LaunchKrnl4(size_t NumElms, int InitVal) {
-  int *Hmm = NULL, *Dptr = NULL, blockSize = 64, DataMismatch = 0;
-  hipStream_t strm;
-  HIP_CHECK(hipStreamCreate(&strm));
-  HIP_CHECK(hipMallocManaged(&Hmm, (sizeof(int) * NumElms)));
-  HIP_CHECK(hipMalloc(&Dptr, (sizeof(int) * NumElms)));
-  int *Hstptr = reinterpret_cast<int*>(new int[NumElms]);
-  for (size_t i = 0; i < NumElms; ++i) {
-    Hstptr[i] = InitVal;
-  }
-  HIP_CHECK(hipMemcpy(Dptr, Hstptr, (NumElms * sizeof(int)),
-                      hipMemcpyHostToDevice));
-  dim3 dimBlock(blockSize, 1, 1);
-  dim3 dimGrid((NumElms + blockSize -1)/blockSize, 1, 1);
-  KrnlWth2MemTypes<<<dimGrid, dimBlock, 0, strm>>>(Hmm, Dptr, NumElms);
-  HIP_CHECK(hipStreamSynchronize(strm));
-  for (size_t i = 0; i < NumElms; ++i) {
-    if (Hmm[i] != (InitVal + 10)) {
-      DataMismatch++;
-    }
-  }
-  if (DataMismatch != 0) {
-    INFO("Data Mismatch observed after the Kernel: KrnlWth2MemTypes!!\n");
-    REQUIRE(false);
-  }
-  DataMismatch = 0;
-  KernelMul_MngdMem<<<dimGrid, dimBlock, 0, strm>>>(Hmm, Dptr, NumElms);
-  HIP_CHECK(hipStreamSynchronize(strm));
-  // Verifying the result
-  for (size_t i = 0; i < NumElms; ++i) {
-    if (Hmm[i] != (InitVal * 10)) {
-      DataMismatch++;
-    }
-  }
-  if (DataMismatch != 0) {
-    INFO("Data Mismatch observedafter the Kernel: KernelMul_MngdMem!!\n");
-    REQUIRE(false);
-  }
-  DataMismatch = 0;
-  KernelMulAdd_MngdMem<<<dimGrid, dimBlock, 0, strm>>>(Hmm, NumElms);
-  HIP_CHECK(hipStreamSynchronize(strm));
-  // Verifying the result
-
-  for (size_t i = 0; i < NumElms; ++i) {
-    if (Hmm[i] != (InitVal * 10 * 2 + 10)) {
-      DataMismatch++;
-    }
-  }
-  if (DataMismatch != 0) {
-    INFO("Data Mismatch observedafter the Kernel: KernelMul_MngdMem!!\n");
-    REQUIRE(false);
-  }
-  delete[] Hstptr;
-}
-
-
 
 static size_t N{4 * 1024 * 1024};
 static unsigned blocksPerCU{6};
@@ -241,104 +175,3 @@ TEST_CASE("Unit_hipMallocManaged_Advanced") {
   }
 }
 
-
-// The following test case tests the behavior of kernel with a HMM memory and
-// hipMalloc memory
-
-TEST_CASE("Unit_hipMallocManaged_KrnlWth2MemTypes") {
-  IfTestPassed = true;
-  int *Hmm = NULL, *Dptr = NULL, InitVal = 123;
-  size_t NumElms = (1024 * 1024);
-  int *Hptr = new int[NumElms], blockSize = 64, DataMismatch = 0;
-  int managed =  HmmAttrPrint();
-  if (managed == 1) {
-    hipStream_t strm;
-    HIP_CHECK(hipStreamCreate(&strm));
-    HIP_CHECK(hipMallocManaged(&Hmm, sizeof(int) * NumElms));
-    HIP_CHECK(hipMalloc(&Dptr, sizeof(int) * NumElms));
-    for (size_t i = 0; i < NumElms; ++i) {
-      Hmm[i] = 0;
-      Hptr[i] = InitVal;
-    }
-    HIP_CHECK(hipMemcpy(Dptr, Hptr, sizeof(int) * NumElms,
-                        hipMemcpyHostToDevice));
-    dim3 dimBlock(blockSize, 1, 1);
-    dim3 dimGrid((NumElms + blockSize -1)/blockSize, 1, 1);
-    KrnlWth2MemTypes<<<dimGrid, dimBlock, 0, strm>>>(Hmm, Dptr, NumElms);
-    HIP_CHECK(hipStreamSynchronize(strm));
-    // Verifying the results
-    for (size_t k = 0; k < NumElms; ++k) {
-      if (Hmm[k] != (InitVal + 10)) {
-        DataMismatch++;
-      }
-    }
-    if (DataMismatch != 0) {
-      WARN("DataMismatch observed!\n");
-      IfTestPassed = false;
-    }
-
-    HIP_CHECK(hipFree(Hmm));
-    HIP_CHECK(hipFree(Dptr));
-    delete[] Hptr;
-    REQUIRE(IfTestPassed);
-  } else {
-    SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
-           "attribute. Hence skipping the testing with Pass result.\n");
-  }
-}
-
-// The following test case tests when the same Hmm memory is used for
-// launching multiple different kernels will results in any issue
-TEST_CASE("Unit_hipMallocManaged_MultiKrnlHmmAccess") {
-  int managed = HmmAttrPrint();
-  if (managed) {
-    int InitVal = 123, NumElms = (1024 * 1024);
-    LaunchKrnl4(NumElms, InitVal);
-  } else {
-    SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
-           "attribute. Hence skipping the testing with Pass result.\n");
-  }
-}
-
-
-//  The following test case allocation, host access, device access of HMM
-//   memory from size 1 to 10KB
-
-TEST_CASE("Unit_hipMallocManaged_MultiSize") {
-  IfTestPassed = true;
-  int managed = HmmAttrPrint();
-  if (managed == 1) {
-    unsigned char *Hmm1 = NULL, *Hmm2 = NULL;
-    int InitVal = 100, blockSize = 64, DataMismatch = 0;
-    hipStream_t strm;
-    HIP_CHECK(hipStreamCreate(&strm));
-    dim3 dimBlock(blockSize, 1, 1);
-    for (int i = 1; i < (1024*1024); ++i) {
-      HIP_CHECK(hipMallocManaged(&Hmm1, i));
-      HIP_CHECK(hipMallocManaged(&Hmm2, i));
-      for (int j = 0; j < i; ++j) {
-        Hmm1[j] = InitVal;
-      }
-      dim3 dimGrid((i + blockSize -1)/blockSize, 1, 1);
-      KrnlWth2MemTypesC<<<dimGrid, dimBlock, 0, strm>>>(Hmm2, Hmm1, i);
-      HIP_CHECK(hipStreamSynchronize(strm));
-      //  Verifying the results
-      for (int k = 0; k < i; ++k) {
-        if (Hmm2[k] != (InitVal + 10)) {
-          DataMismatch++;
-        }
-      }
-      if (DataMismatch != 0) {
-        WARN("DataMismatch observed!\n");
-        IfTestPassed = false;
-      }
-      DataMismatch = 0;
-      HIP_CHECK(hipFree(Hmm1));
-      HIP_CHECK(hipFree(Hmm2));
-      REQUIRE(IfTestPassed);
-    }
-  } else {
-    SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
-           "attribute. Hence skipping the testing with Pass result.\n");
-  }
-}

tests/catch/unit/memory/hipMallocManagedFlagsTst.cc

@@ -262,79 +262,3 @@ TEST_CASE("Unit_hipMallocManaged_AccessMultiStream") {
   }
 }
 
-TEST_CASE("Unit_hipMallocManaged_ExtremeSizes") {
-  int managed = HmmAttrPrint();
-  if (managed == 1) {
-    bool IfTestPassed = true;
-    hipError_t err;
-    void *Hmm = NULL;
-    size_t totalDevMem = 0, freeDevMem = 0;
-    int NumDevs = 0;
-    HIP_CHECK(hipGetDeviceCount(&NumDevs));
-    // Testing allocation of extreme and unusual mem values
-    for (int i = 0; i < NumDevs; i++) {
-      HIP_CHECK(hipSetDevice(i));
-      HIP_CHECK(hipMemGetInfo(&freeDevMem, &totalDevMem));
-      err = hipMallocManaged(&Hmm, 1, hipMemAttachGlobal);
-      if (hipSuccess == err) {
-        HIP_CHECK(hipFree(Hmm));
-      } else {
-        WARN("Observed error while allocating memory on GPU: " << i);
-        WARN(" size 1 with");
-        WARN(" hipMallocManaged() api with flag 'hipMemAttachGlobal'\n");
-        WARN("Error: " << hipGetErrorString(err));
-        IfTestPassed = false;
-      }
-      err = hipMallocManaged(&Hmm, freeDevMem, hipMemAttachGlobal);
-      if (hipSuccess == err) {
-        HIP_CHECK(hipFree(Hmm));
-      } else {
-        WARN("Observed error while allocating max free memory on GPU: " << i);
-        WARN(" with hipMallocManaged() api with flag 'hipMemAttachGlobal'\n");
-        WARN("Error: " << hipGetErrorString(err));
-        IfTestPassed = false;
-      }
-      err = hipMallocManaged(&Hmm, (freeDevMem - 1), hipMemAttachGlobal);
-      if (hipSuccess == err) {
-        HIP_CHECK(hipFree(Hmm));
-      } else {
-        WARN("Observed error while allocating max (free - 1) memory on ");
-        WARN("GPU: " << i);
-        WARN(" using hipMallocManaged() api with flag 'hipMemAttachGlobal'\n");
-        WARN("Error: " << hipGetErrorString(err));
-        IfTestPassed = false;
-      }
-      err = hipMallocManaged(&Hmm, 1, hipMemAttachHost);
-      if (hipSuccess == err) {
-        HIP_CHECK(hipFree(Hmm));
-      } else {
-        WARN("Observed error while allocating memory size 1 on GPU: " << i);
-        WARN(" with hipMallocManaged() api with flag 'hipMemAttachHost'\n");
-        WARN("Error: " << hipGetErrorString(err));
-        IfTestPassed = false;
-      }
-      err = hipMallocManaged(&Hmm, freeDevMem, hipMemAttachHost);
-      if (hipSuccess == err) {
-        HIP_CHECK(hipFree(Hmm));
-      } else {
-        WARN("Observed error while allocating max free memory on GPU: " << i);
-        WARN(" with hipMallocManaged() api with flag 'hipMemAttachHost'\n");
-        WARN("Error: " << hipGetErrorString(err));
-        IfTestPassed = false;
-      }
-      err = hipMallocManaged(&Hmm, (freeDevMem - 1), hipMemAttachHost);
-      if (hipSuccess == err) {
-        HIP_CHECK(hipFree(Hmm));
-      } else {
-        WARN("Observed error while allocating max (freeDevMem - 1) memory"
-               " on GPU: " << i);
-        WARN(" with hipMallocManaged() api with flag 'hipMemAttachHost'\n");
-        WARN("Error: " << hipGetErrorString(err));
-        IfTestPassed = false;
-      }
-    }
-    REQUIRE(IfTestPassed);
-  } else {
-    SUCCEED("Gpu doesnt support HMM! Hence skipping the test with PASS result");
-  }
-}

tests/catch/unit/memory/hipMemRangeGetAttribute.cc

@@ -0,0 +1,408 @@
+/*
+Copyright (c) 2021 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, INCLUDING 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 ANY 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.
+*/
+
+/* Test Case Description:
+   Scenario-1: The following function tests the count parameter(last param) to
+   hipMemRangeGetAttribute api by passing possible extreme values.
+   Curently the only way to test if count param working properly is to verify
+   the first parameter of hipMemRangeGetAttribute() api has value 1 stored
+
+   Scenario-2: This test case checks the behavior of hipMemRangeGetAttribute()  with
+   AccessedBy flag is consistent with cuda's counter part
+
+   Scenario-3: Allocate  4 * page size of memory with the flag hipMemAttachGloal. Advise
+   AccessedBy, ReadMostly and PreferredLocation to first half(2*pageSz) of the
+   memory and probe the for the flags which are set earlier using
+   hipMemRangeGetAttribute() api for the full size(4*PageSz).
+
+
+   Scenario-4: The following scenarios tests that probing the attributes which are not set
+   by hipMemAdvise() but being probed using hipMemRangeGetAttribute() should
+   not result in a crash
+
+   Scenario-5: The following scenario is a simple test which does the following:
+   Allocate Hmm memory --> hipMemPrefetchAsync() to device 0 and then
+   probe LastPrefetchLocation attribute using hipMemRangeGetAttribute
+
+   Scenario-6: The following Test Case does negative tests on hipMemRangeGetAttribute()*/
+
+#include <hip_test_common.hh>
+#include <stdlib.h>
+#ifdef __linux__
+  #include <unistd.h>
+  #include <sys/sysinfo.h>
+#endif
+
+static bool CheckError(hipError_t err, int LineNo) {
+  if (err == hipSuccess) {
+    WARN("Error expected but received hipSuccess at line no.:"
+           << LineNo);
+    return false;
+  } else {
+    return true;
+  }
+}
+
+
+static int HmmAttrPrint() {
+  int managed = 0;
+  WARN("The following are the attribute values related to HMM for"
+         " device 0:\n");
+  HIP_CHECK(hipDeviceGetAttribute(&managed,
+              hipDeviceAttributeDirectManagedMemAccessFromHost, 0));
+  WARN("hipDeviceAttributeDirectManagedMemAccessFromHost: " << managed);
+  HIP_CHECK(hipDeviceGetAttribute(&managed,
+                                 hipDeviceAttributeConcurrentManagedAccess, 0));
+  WARN("hipDeviceAttributeConcurrentManagedAccess: " << managed);
+  HIP_CHECK(hipDeviceGetAttribute(&managed,
+                                 hipDeviceAttributePageableMemoryAccess, 0));
+  WARN("hipDeviceAttributePageableMemoryAccess: " << managed);
+  HIP_CHECK(hipDeviceGetAttribute(&managed,
+              hipDeviceAttributePageableMemoryAccessUsesHostPageTables, 0));
+  WARN("hipDeviceAttributePageableMemoryAccessUsesHostPageTables:"
+         << managed);
+
+  HIP_CHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeManagedMemory,
+                                  0));
+  WARN("hipDeviceAttributeManagedMemory: " << managed);
+  return managed;
+}
+
+// The following function tests the count parameter(last param) to
+// hipMemRangeGetAttribute api by passing possible extreme values.
+// Curently the only way to test if count param working properly is to verify
+// the first parameter of hipMemRangeGetAttribute() api has value 1 stored
+TEST_CASE("Unit_hipMemRangeGetAttribute_TstCountParam") {
+  int MangdMem = HmmAttrPrint();
+  if (MangdMem == 1) {
+    int MEM_SIZE = 4096, RND_NUM = 9999, FLG_READMOSTLY_ENBLD = 1;
+    bool IfTestPassed = true;
+    int data = RND_NUM, *devPtr = nullptr;
+    size_t TotGpuMem, TotGpuFreeMem;
+    HIP_CHECK(hipMemGetInfo(&TotGpuFreeMem, &TotGpuMem));
+
+    HIP_CHECK(hipMallocManaged(&devPtr, MEM_SIZE, hipMemAttachGlobal));
+    HIP_CHECK(hipMemAdvise(devPtr, MEM_SIZE, hipMemAdviseSetReadMostly, 0));
+    HIP_CHECK(hipMemRangeGetAttribute(reinterpret_cast<void*>(&data),
+                                     sizeof(int),
+                                     hipMemRangeAttributeReadMostly,
+                                     devPtr, MEM_SIZE));
+    if (data != FLG_READMOSTLY_ENBLD) {
+      WARN("hipMemRangeGetAttribute() api didnt return expected value!\n");
+      IfTestPassed = false;
+    }
+    HIP_CHECK(hipFree(devPtr));
+    HIP_CHECK(hipMallocManaged(&devPtr, TotGpuFreeMem, hipMemAttachGlobal));
+    HIP_CHECK(hipMemAdvise(devPtr, TotGpuFreeMem, hipMemAdviseSetReadMostly,
+                           0));
+    HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int),
+                                     hipMemRangeAttributeReadMostly,
+                                     devPtr, TotGpuFreeMem));
+
+    if (data != FLG_READMOSTLY_ENBLD) {
+      WARN("hipMemRangeGetAttribute() api didnt return expected value!\n");
+      IfTestPassed = false;
+    }
+    HIP_CHECK(hipFree(devPtr));
+    HIP_CHECK(hipMallocManaged(&devPtr, (TotGpuFreeMem - 1),
+                              hipMemAttachGlobal));
+    HIP_CHECK(hipMemAdvise(devPtr, (TotGpuFreeMem - 1),
+                          hipMemAdviseSetReadMostly, 0));
+    HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int),
+                                     hipMemRangeAttributeReadMostly,
+                                     devPtr, (TotGpuFreeMem - 1)));
+
+    if (data != FLG_READMOSTLY_ENBLD) {
+      WARN("hipMemRangeGetAttribute() api didnt return expected value!\n");
+      IfTestPassed = false;
+    }
+    HIP_CHECK(hipFree(devPtr));
+
+    REQUIRE(IfTestPassed);
+  } else {
+    SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
+           "attribute. Hence skipping the testing with Pass result.\n");
+  }
+}
+
+/* The following Test Case does negative tests on hipMemRangeGetAttribute()*/
+
+TEST_CASE("Unit_hipMemRangeGetAttribute_NegativeTests") {
+  int MangdMem = HmmAttrPrint();
+  if (MangdMem == 1) {
+    int MEM_SIZE = 4096, RND_NUM = 9999;
+    float *devPtr = nullptr;
+    int NumDevs;
+    HIP_CHECK(hipGetDeviceCount(&NumDevs));
+    int data = RND_NUM;
+    int *OutData = new int[NumDevs];
+    for (int m = 0; m < NumDevs; ++m) {
+      OutData[m] = RND_NUM;
+    }
+    HIP_CHECK(hipMallocManaged(&devPtr, MEM_SIZE, hipMemAttachGlobal));
+    HIP_CHECK(hipMemAdvise(devPtr, MEM_SIZE, hipMemAdviseSetReadMostly, 0));
+
+    // checking the behavior with dataSize 0
+    SECTION("checking the behavior with dataSize 0") {
+      REQUIRE(CheckError(hipMemRangeGetAttribute(&data, 0,
+                                     hipMemRangeAttributeReadMostly,
+                                     devPtr, MEM_SIZE), __LINE__));
+    }
+    // checking the behavior with dataSize > 4 and even
+    SECTION("checking the behavior with dataSize > 4 and even") {
+      REQUIRE(CheckError(hipMemRangeGetAttribute(OutData, 6,
+                                     hipMemRangeAttributeReadMostly,
+                                     devPtr, MEM_SIZE), __LINE__));
+    }
+    // checking the behavior with dataSize > 4 and odd
+    SECTION("checking the behavior with dataSize > 4 and odd") {
+      REQUIRE(CheckError(hipMemRangeGetAttribute(OutData, 7,
+                                     hipMemRangeAttributeReadMostly,
+                                     devPtr, MEM_SIZE), __LINE__));
+    }
+    // checking the behavior with dataSize which is not multiple of 4
+    SECTION("checking the behavior with dataSize which is not multiple of 4") {
+      REQUIRE(CheckError(hipMemRangeGetAttribute(OutData, 27,
+                                     hipMemRangeAttributeReadMostly,
+                                     devPtr, MEM_SIZE), __LINE__));
+    }
+    // checking the behaviour with devPtr(4th param) as NULL
+    SECTION("checking the behaviour with devPtr(4th param) as NULL") {
+    REQUIRE(CheckError(hipMemRangeGetAttribute(&data, sizeof(int),
+                                     hipMemRangeAttributeReadMostly,
+                                     NULL, MEM_SIZE), __LINE__));
+    }
+    // checking the behaviour with count(5th param) as 0
+    SECTION("checking the behaviour with count(5th param) as 0") {
+      REQUIRE(CheckError(hipMemRangeGetAttribute(&data, sizeof(int),
+                                     hipMemRangeAttributeReadMostly,
+                                     devPtr, 0), __LINE__));
+    }
+    // checking the behavior with invalid attribute (3rd param) as 0
+    // as it is attribute hence avoiding the negative tests with 3rd param
+
+    // checking the behaviour of the api with ptr allocated using
+    // hipHostMalloc
+    void *ptr = nullptr;
+    SECTION("Checking behavior with hipHostMalloc ptr") {
+      HIP_CHECK(hipHostMalloc(&ptr, MEM_SIZE, 0));
+      REQUIRE(CheckError(hipMemRangeGetAttribute(&data, sizeof(int),
+                                     hipMemRangeAttributeReadMostly,
+                                     ptr, MEM_SIZE), __LINE__));
+      HIP_CHECK(hipHostFree(ptr));
+    }
+    HIP_CHECK(hipFree(devPtr));
+  } else {
+    SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
+           "attribute. Hence skipping the testing with Pass result.\n");
+  }
+}
+
+/* This test case checks the behavior of hipMemRangeGetAttribute()  with
+   AccessedBy flag is consistent with cuda's counter part*/
+TEST_CASE("Unit_hipMemRangeGetAttribute_AccessedBy1") {
+  int managed = HmmAttrPrint();
+  if (managed == 1) {
+    int Ngpus = 0, *Hmm = NULL, MEM_SZ = 4096, RND_NUM = 999;
+    HIP_CHECK(hipGetDeviceCount(&Ngpus));
+    int *OutData = new int[Ngpus];
+    for (int i = 0; i < Ngpus; ++i) {
+      OutData[Ngpus] = RND_NUM;
+    }
+    HIP_CHECK(hipMallocManaged(&Hmm, MEM_SZ));
+    HIP_CHECK(hipMemAdvise(Hmm, MEM_SZ, hipMemAdviseSetAccessedBy, 0));
+    HIP_CHECK(hipMemRangeGetAttribute(OutData, 4*Ngpus,
+                                     hipMemRangeAttributeAccessedBy,
+                                     Hmm, MEM_SZ));
+    if (OutData[0] != 0) {
+      WARN("Didn't receive expected value at line: " << __LINE__);
+      REQUIRE(false);
+    }
+    for (int i = 1; i < Ngpus; ++i) {
+      if (OutData[i] != -2) {
+        WARN("Didn't receive expected value at line: " << __LINE__);
+        REQUIRE(false);
+      }
+    }
+    if (Ngpus >= 2) {
+      for (int i = 0; i < Ngpus; ++i) {
+        HIP_CHECK(hipMemAdvise(Hmm, MEM_SZ, hipMemAdviseSetAccessedBy, i));
+      }
+      // checking the behavior with dataSize less than the number of gpus
+      // This should not result in segfault.
+      HIP_CHECK(hipMemRangeGetAttribute(OutData, 4*(Ngpus-1),
+                                       hipMemRangeAttributeAccessedBy,
+                                       Hmm, MEM_SZ));
+      // OutData should have stored the gpu ordinals for which AccessedBy is
+      // assigned except for the last element which should have -2 stored
+      // so as to be consistent with cuda's behavior
+      for (int i = 0; i < (Ngpus - 1); ++i) {
+        if (OutData[i] != i) {
+          WARN("Didn't receive expected value at line: " << __LINE__);
+          REQUIRE(false);
+        }
+      }
+      if (OutData[Ngpus - 1] != -2) {
+        WARN("Didn't receive expected value at line: " << __LINE__);
+        REQUIRE(false);
+      }
+    }
+    HIP_CHECK(hipFree(Hmm));
+    delete[] OutData;
+  } else {
+    SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
+           "attribute. Hence skipping the testing with Pass result.\n");
+  }
+}
+
+
+
+
+/* Allocate  4 * page size of memory with the flag hipMemAttachGloal. Advise
+   AccessedBy, ReadMostly and PreferredLocation to first half(2*pageSz) of the
+    memory and probe the for the flags which are set earlier using
+   hipMemRangeGetAttribute() api for the full size(4*PageSz).*/
+/* Need to discuss the difference in behavior w.r.t cuda*/
+
+TEST_CASE("Unit_hipMemRangeGetAttribte_3") {
+  int managed = HmmAttrPrint();
+  if (managed == 1) {
+    int Ngpus = 0, *Hmm = NULL, MEM_SZ = 4096*4, RND_NUM = 999;
+    HIP_CHECK(hipGetDeviceCount(&Ngpus));
+    int *OutData = new int[Ngpus];
+    for (int i = 0; i < Ngpus; ++i) {
+      OutData[Ngpus] = RND_NUM;
+    }
+    HIP_CHECK(hipMallocManaged(&Hmm, MEM_SZ));
+    HIP_CHECK(hipMemAdvise(Hmm, MEM_SZ/2, hipMemAdviseSetAccessedBy, 0));
+    HIP_CHECK(hipMemRangeGetAttribute(OutData, 4*Ngpus,
+                                     hipMemRangeAttributeAccessedBy,
+                                     (Hmm), MEM_SZ));
+
+    HIP_CHECK(hipMemAdvise(Hmm, MEM_SZ/2, hipMemAdviseSetReadMostly, 0));
+    // The Api called below should not fail
+    HIP_CHECK(hipMemRangeGetAttribute(OutData, 4,
+                                     hipMemRangeAttributeReadMostly,
+                                     (Hmm), MEM_SZ));
+
+    HIP_CHECK(hipMemAdvise(Hmm, MEM_SZ/2, hipMemAdviseSetPreferredLocation, 0));
+    // The api called below should not fail
+    HIP_CHECK(hipMemRangeGetAttribute(OutData, 4,
+                                     hipMemRangeAttributePreferredLocation,
+                                     (Hmm), MEM_SZ));
+    HIP_CHECK(hipFree(Hmm));
+    delete[] OutData;
+  } else {
+    SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
+           "attribute. Hence skipping the testing with Pass result.\n");
+  }
+}
+
+
+/* The following scenarios tests that probing the attributes which are not set
+   by hipMemAdvise() but being probed using hipMemRangeGetAttribute() should
+   not result in a crash*/
+
+TEST_CASE("Unit_hipMemRangeGetAttribute_4") {
+  int managed = HmmAttrPrint();
+  if (managed == 1) {
+    int *Hmm = NULL, PageSz = 4096, Ngpus, RND_NUM = 999;
+    HIP_CHECK(hipGetDeviceCount(&Ngpus));
+    int *OutData = new int[Ngpus];
+    for (int i = 0; i < Ngpus; ++i) {
+      OutData[i] = RND_NUM;
+    }
+    HIP_CHECK(hipMallocManaged(&Hmm, 4*PageSz));
+    SECTION("Set ReadMostly & probe other flags") {
+      HIP_CHECK(hipMemAdvise(Hmm, 4*PageSz, hipMemAdviseSetReadMostly, 0));
+      HIP_CHECK(hipMemRangeGetAttribute(OutData, 4*Ngpus,
+                                       hipMemRangeAttributeAccessedBy,
+                                       Hmm, 4*PageSz));
+      HIP_CHECK(hipMemRangeGetAttribute(OutData, 4,
+                                       hipMemRangeAttributePreferredLocation,
+                                       Hmm, 4*PageSz));
+      HIP_CHECK(hipMemAdvise(Hmm, 4*PageSz, hipMemAdviseUnsetReadMostly, 0));
+    }
+    SECTION("Set AccessedBy & probe other flags") {
+      HIP_CHECK(hipMemAdvise(Hmm, 4*PageSz, hipMemAdviseSetAccessedBy, 0));
+      HIP_CHECK(hipMemRangeGetAttribute(OutData, 4,
+                                       hipMemRangeAttributeReadMostly,
+                                       Hmm, 4*PageSz));
+      HIP_CHECK(hipMemRangeGetAttribute(OutData, 4,
+                                       hipMemRangeAttributePreferredLocation,
+                                       Hmm, 4*PageSz));
+      HIP_CHECK(hipMemAdvise(Hmm, 4*PageSz, hipMemAdviseUnsetAccessedBy, 0));
+    }
+    SECTION("Set AccessedBy & probe other flags") {
+      HIP_CHECK(hipMemAdvise(Hmm, 4*PageSz, hipMemAdviseSetPreferredLocation,
+                            0));
+      HIP_CHECK(hipMemRangeGetAttribute(OutData, 4,
+                                       hipMemRangeAttributeReadMostly,
+                                       Hmm, 4*PageSz));
+      HIP_CHECK(hipMemRangeGetAttribute(OutData, 4*Ngpus,
+                                       hipMemRangeAttributeAccessedBy,
+                                       Hmm, 4*PageSz));
+      HIP_CHECK(hipMemAdvise(Hmm, 4*PageSz, hipMemAdviseUnsetPreferredLocation,
+                            0));
+    }
+    HIP_CHECK(hipFree(Hmm));
+    delete[] OutData;
+  } else {
+    SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
+           "attribute. Hence skipping the testing with Pass result.\n");
+  }
+}
+
+
+/* The following scenario is a simple test which does the following:
+   Allocate Hmm memory --> hipMemPrefetchAsync() to device 0 and then
+   probe LastPrefetchLocation attribute using hipMemRangeGetAttribute*/
+
+TEST_CASE("Unit_hipMemRangeGetAttribute_PrefetchAndGtAttr") {
+  int managed = HmmAttrPrint();
+  if (managed == 1) {
+    int Ngpus = 0, *Hmm = NULL, RND_NUM = 999;
+    size_t PageSz = 4096;
+    HIP_CHECK(hipGetDeviceCount(&Ngpus));
+
+    int *OutData = new int[Ngpus];
+    for (int i = 0; i < Ngpus; ++i) {
+      OutData[Ngpus] = RND_NUM;
+    }
+    HIP_CHECK(hipMallocManaged(&Hmm, PageSz*4));
+    hipStream_t strm;
+    HIP_CHECK(hipStreamCreate(&strm));
+    HIP_CHECK(hipMemPrefetchAsync(Hmm, PageSz*4, 0, strm));
+    HIP_CHECK(hipStreamSynchronize(strm));
+    HIP_CHECK(hipMemRangeGetAttribute(OutData, 4,
+                                     hipMemRangeAttributeLastPrefetchLocation,
+                                     Hmm, PageSz*4));
+    HIP_CHECK(hipStreamDestroy(strm));
+    HIP_CHECK(hipFree(Hmm));
+    if (OutData[0] != 0) {
+      WARN("Didnt receive expected value at line: " << __LINE__);
+      delete[] OutData;
+      REQUIRE(false);
+    }
+  } else {
+    SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
+           "attribute. Hence skipping the testing with Pass result.\n");
+  }
+}
+

tests/catch/unit/memory/hipMemcpy2DAsync.cc

@@ -83,22 +83,40 @@ TEMPLATE_TEST_CASE("Unit_hipMemcpy2DAsync_Host&PinnedMem", ""
 
   // Initialize the data
   HipTest::setDefaultData<TestType>(NUM_W*NUM_H, A_h, B_h, C_h);
+  SECTION("Calling Async apis with stream object created by user") {
+    // Host to Device
+    HIP_CHECK(hipMemcpy2DAsync(A_d, pitch_A, A_h, COLUMNS*sizeof(TestType),
+                               COLUMNS*sizeof(TestType), ROWS,
+                               hipMemcpyHostToDevice, stream));
 
-  // Host to Device
-  HIP_CHECK(hipMemcpy2DAsync(A_d, pitch_A, A_h, COLUMNS*sizeof(TestType),
-                             COLUMNS*sizeof(TestType), ROWS,
-                             hipMemcpyHostToDevice, stream));
+    // Performs D2D on same GPU device
+    HIP_CHECK(hipMemcpy2DAsync(B_d, pitch_B, A_d,
+                               pitch_A, COLUMNS*sizeof(TestType),
+                               ROWS, hipMemcpyDeviceToDevice, stream));
 
-  // Performs D2D on same GPU device
-  HIP_CHECK(hipMemcpy2DAsync(B_d, pitch_B, A_d,
-                             pitch_A, COLUMNS*sizeof(TestType),
-                             ROWS, hipMemcpyDeviceToDevice, stream));
+    // hipMemcpy2DAsync Device to Host
+    HIP_CHECK(hipMemcpy2DAsync(B_h, COLUMNS*sizeof(TestType), B_d, pitch_B,
+                               COLUMNS*sizeof(TestType), ROWS,
+                               hipMemcpyDeviceToHost, stream));
+    HIP_CHECK(hipStreamSynchronize(stream));
+  }
+  SECTION("Calling Async apis with hipStreamPerThread") {
+    // Host to Device
+    HIP_CHECK(hipMemcpy2DAsync(A_d, pitch_A, A_h, COLUMNS*sizeof(TestType),
+                               COLUMNS*sizeof(TestType), ROWS,
+                               hipMemcpyHostToDevice, hipStreamPerThread));
 
-  // hipMemcpy2DAsync Device to Host
-  HIP_CHECK(hipMemcpy2DAsync(B_h, COLUMNS*sizeof(TestType), B_d, pitch_B,
-                             COLUMNS*sizeof(TestType), ROWS,
-                             hipMemcpyDeviceToHost, stream));
-  HIP_CHECK(hipStreamSynchronize(stream));
+    // Performs D2D on same GPU device
+    HIP_CHECK(hipMemcpy2DAsync(B_d, pitch_B, A_d,
+                               pitch_A, COLUMNS*sizeof(TestType),
+                               ROWS, hipMemcpyDeviceToDevice, hipStreamPerThread));
+
+    // hipMemcpy2DAsync Device to Host
+    HIP_CHECK(hipMemcpy2DAsync(B_h, COLUMNS*sizeof(TestType), B_d, pitch_B,
+                               COLUMNS*sizeof(TestType), ROWS,
+                               hipMemcpyDeviceToHost, hipStreamPerThread));
+    HIP_CHECK(hipStreamSynchronize(hipStreamPerThread));
+  }
 
   // Validating the result
   REQUIRE(HipTest::checkArray<TestType>(A_h, B_h, COLUMNS, ROWS) == true);

tests/catch/unit/memory/hipMemcpy2DFromArrayAsync.cc

@@ -62,10 +62,18 @@ TEST_CASE("Unit_hipMemcpy2DFromArrayAsync_Basic") {
   HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, hData, width,
                                     width, NUM_H,
                                     hipMemcpyHostToDevice));
-  HIP_CHECK(hipMemcpy2DFromArrayAsync(A_h, width, A_d,
-                                 0, 0, width, NUM_H,
-                                 hipMemcpyDeviceToHost, stream));
-  HIP_CHECK(hipStreamSynchronize(stream));
+  SECTION("Calling hipMemcpy2DFromArrayAsync() with user declared stream obj") {
+    HIP_CHECK(hipMemcpy2DFromArrayAsync(A_h, width, A_d,
+                                   0, 0, width, NUM_H,
+                                   hipMemcpyDeviceToHost, stream));
+    HIP_CHECK(hipStreamSynchronize(stream));
+  }
+  SECTION("Calling hipMemcpy2DFromArrayAsync() with hipStreamPerThread") {
+    HIP_CHECK(hipMemcpy2DFromArrayAsync(A_h, width, A_d,
+                                   0, 0, width, NUM_H,
+                                   hipMemcpyDeviceToHost, hipStreamPerThread));
+    HIP_CHECK(hipStreamSynchronize(hipStreamPerThread));
+  }
   REQUIRE(HipTest::checkArray(A_h, hData, NUM_W, NUM_H) == true);
 
   // Cleaning the memory

tests/catch/unit/memory/hipMemcpy2DToArrayAsync.cc

@@ -58,11 +58,18 @@ TEST_CASE("Unit_hipMemcpy2DToArrayAsync_Basic") {
   HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault));
   HipTest::setDefaultData<float>(width*NUM_H, A_h, hData, nullptr);
   HIP_CHECK(hipStreamCreate(&stream));
-
-  HIP_CHECK(hipMemcpy2DToArrayAsync(A_d, 0, 0, hData, width,
-                                    width, NUM_H,
-                                    hipMemcpyHostToDevice, stream));
-  HIP_CHECK(hipStreamSynchronize(stream));
+  SECTION("Calling hipMemcpy2DToArrayAsync() with user declared stream obj") {
+    HIP_CHECK(hipMemcpy2DToArrayAsync(A_d, 0, 0, hData, width,
+                                      width, NUM_H,
+                                      hipMemcpyHostToDevice, stream));
+    HIP_CHECK(hipStreamSynchronize(stream));
+  }
+  SECTION("Calling hipMemcpy2DToArrayAsync() with hipStreamPerThread") {
+    HIP_CHECK(hipMemcpy2DToArrayAsync(A_d, 0, 0, hData, width,
+                                      width, NUM_H,
+                                      hipMemcpyHostToDevice, hipStreamPerThread));
+    HIP_CHECK(hipStreamSynchronize(hipStreamPerThread));
+  }
   HIP_CHECK(hipMemcpy2DFromArray(A_h, width, A_d,
                                  0, 0, width, NUM_H,
                                  hipMemcpyDeviceToHost));

tests/catch/unit/memory/hipMemcpy3DAsync.cc

@@ -602,8 +602,14 @@ void Memcpy3DAsync<T>::simple_Memcpy3DAsync() {
 #else
   myparms.kind = hipMemcpyHostToDevice;
 #endif
-  REQUIRE(hipMemcpy3DAsync(&myparms, stream) == hipSuccess);
-  HIP_CHECK(hipStreamSynchronize(stream));
+  SECTION("Calling hipMemcpy3DAsync() using user declared stream obj") {
+    REQUIRE(hipMemcpy3DAsync(&myparms, stream) == hipSuccess);
+    HIP_CHECK(hipStreamSynchronize(stream));
+  }
+  SECTION("Calling hipMemcpy3DAsync() using hipStreamPerThread") {
+    REQUIRE(hipMemcpy3DAsync(&myparms, hipStreamPerThread) == hipSuccess);
+    HIP_CHECK(hipStreamSynchronize(hipStreamPerThread));
+  }
 
   // Array to Array
   memset(&myparms, 0x0, sizeof(hipMemcpy3DParms));

tests/catch/unit/memory/hipMemcpyPeerAsync.cc

@@ -149,11 +149,20 @@ TEST_CASE("Unit_hipMemcpyPeerAsync_Basic") {
 
       // Copying data from GPU-0 to GPU-1 and performing vector addition
       HIP_CHECK(hipSetDevice(1));
-      HIP_CHECK(hipMemcpyPeerAsync(X_d, 1, A_d, 0, copy_bytes,
-                                   stream));
-      HIP_CHECK(hipMemcpyPeerAsync(Y_d, 1, B_d, 0, copy_bytes,
-                                   stream));
-      HIP_CHECK(hipStreamSynchronize(stream));
+      SECTION("Calling hipMemcpyPerAsync() using user defined stream obj") {
+        HIP_CHECK(hipMemcpyPeerAsync(X_d, 1, A_d, 0, copy_bytes,
+                                     stream));
+        HIP_CHECK(hipMemcpyPeerAsync(Y_d, 1, B_d, 0, copy_bytes,
+                                     stream));
+	HIP_CHECK(hipStreamSynchronize(stream));
+      }
+      SECTION("Calling hipMemcpyPerAsync() using hipStreamPerThread") {
+        HIP_CHECK(hipMemcpyPeerAsync(X_d, 1, A_d, 0, copy_bytes,
+                                     hipStreamPerThread));
+        HIP_CHECK(hipMemcpyPeerAsync(Y_d, 1, B_d, 0, copy_bytes,
+                                     hipStreamPerThread));
+	HIP_CHECK(hipStreamSynchronize(hipStreamPerThread));
+      }
       hipLaunchKernelGGL(HipTest::vectorADD, dim3(1), dim3(1),
           0, 0, static_cast<const int*>(X_d),
           static_cast<const int*>(Y_d), Z_d, numElements*sizeof(int));

tests/catch/unit/memory/hipMemset2DAsyncMultiThreadAndKernel.cc

@@ -61,7 +61,6 @@ TEST_CASE("Unit_hipMemset2DAsync_WithKernel") {
   size_t elements = NUM_W * NUM_H;
   unsigned blocks{};
   int validateCount{};
-  hipStream_t stream;
 
   blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
   HIP_CHECK(hipMallocPitch(reinterpret_cast<void**>(&A_d), &pitch_A,
@@ -81,21 +80,42 @@ TEST_CASE("Unit_hipMemset2DAsync_WithKernel") {
   }
   HIP_CHECK(hipMemcpy2D(B_d, width, B_h, pitch_B, NUM_W, NUM_H,
                        hipMemcpyHostToDevice));
-  HIP_CHECK(hipStreamCreate(&stream));
+  SECTION("Using User created stream") {
+    hipStream_t stream;
+    HIP_CHECK(hipStreamCreate(&stream));
+    for (size_t k = 0; k < ITER; k++) {
+      hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks),
+                     dim3(threadsPerBlock), 0, stream, B_d, C_d, elements);
+      HIP_CHECK(hipStreamSynchronize(stream));
+      HIP_CHECK(hipMemset2DAsync(C_d, pitch_C, memsetval, NUM_W, NUM_H,
+                stream));
+      HIP_CHECK(hipStreamSynchronize(stream)); 
+      HIP_CHECK(hipMemcpy2D(A_h, width, C_d, pitch_C, NUM_W, NUM_H,
+                           hipMemcpyDeviceToHost));
 
+      for (size_t p = 0 ; p < elements ; p++) {
+        if (A_h[p] == memsetval) {
+          validateCount+= 1;
+        }
+      }
+    }
+    HIP_CHECK(hipStreamDestroy(stream));
+  }
+  SECTION("Using hipStreamPerThread stream") {
+    for (size_t k = 0; k < ITER; k++) {
+      hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks),
+                   dim3(threadsPerBlock), 0, hipStreamPerThread, B_d, C_d, elements);
+      HIP_CHECK(hipStreamSynchronize(hipStreamPerThread));
+      HIP_CHECK(hipMemset2DAsync(C_d, pitch_C, memsetval, NUM_W, NUM_H,
+                hipStreamPerThread));
+      HIP_CHECK(hipStreamSynchronize(hipStreamPerThread)); 
+      HIP_CHECK(hipMemcpy2D(A_h, width, C_d, pitch_C, NUM_W, NUM_H,
+                           hipMemcpyDeviceToHost));
 
-  for (size_t k = 0; k < ITER; k++) {
-    hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks),
-                   dim3(threadsPerBlock), 0, stream, B_d, C_d, elements);
-
-    HIP_CHECK(hipMemset2DAsync(C_d, pitch_C, memsetval, NUM_W, NUM_H, stream));
-    HIP_CHECK(hipStreamSynchronize(stream));
-    HIP_CHECK(hipMemcpy2D(A_h, width, C_d, pitch_C, NUM_W, NUM_H,
-                         hipMemcpyDeviceToHost));
-
-    for (size_t p = 0 ; p < elements ; p++) {
-      if (A_h[p] == memsetval) {
-        validateCount+= 1;
+      for (size_t p = 0 ; p < elements ; p++) {
+        if (A_h[p] == memsetval) {
+          validateCount+= 1;
+        }
       }
     }
   }
@@ -104,7 +124,6 @@ TEST_CASE("Unit_hipMemset2DAsync_WithKernel") {
 
   HIP_CHECK(hipFree(A_d)); HIP_CHECK(hipFree(B_d)); HIP_CHECK(hipFree(C_d));
   free(A_h); free(B_h);
-  HIP_CHECK(hipStreamDestroy(stream));
 }
 
 

tests/catch/unit/memory/hipMemset3DFunctional.cc

@@ -184,11 +184,17 @@ static void testMemsetMaxValue(bool bAsync) {
 
   HIP_CHECK(hipMalloc3D(&devPitchedPtr, extent));
   if (bAsync) {
-    hipStream_t stream;
-    HIP_CHECK(hipStreamCreate(&stream));
-    HIP_CHECK(hipMemset3DAsync(devPitchedPtr, memsetval, extent, stream));
-    HIP_CHECK(hipStreamSynchronize(stream));
-    HIP_CHECK(hipStreamDestroy(stream));
+    SECTION("Using user created stream") {
+      hipStream_t stream;
+      HIP_CHECK(hipStreamCreate(&stream));
+      HIP_CHECK(hipMemset3DAsync(devPitchedPtr, memsetval, extent, stream));
+      HIP_CHECK(hipStreamSynchronize(stream));
+      HIP_CHECK(hipStreamDestroy(stream));
+    }
+    SECTION("Using hipStreamPerThread") {
+      HIP_CHECK(hipMemset3DAsync(devPitchedPtr, memsetval, extent, hipStreamPerThread));
+      HIP_CHECK(hipStreamSynchronize(hipStreamPerThread));
+    }
   } else {
     HIP_CHECK(hipMemset3D(devPitchedPtr, memsetval, extent));
   }
@@ -236,7 +242,7 @@ static void seekAndSet3DArraySlice(bool bAsync) {
 
   // select random slice for memset
   unsigned int seed = time(nullptr);
-  int slice_index = HipTest::RAND_R(&seed) % ZSIZE_S;
+  int slice_index = rand_r(&seed) % ZSIZE_S;
 
   INFO("memset3d for sliceindex " << slice_index);
 

tests/catch/unit/memory/hipMemsetAsyncAndKernel.cc

@@ -21,7 +21,6 @@
  * Test for checking order of execution of device kernel and
  * hipMemsetAsync apis on all gpus
  */
-
 #include <hip_test_common.hh>
 #include <hip_test_checkers.hh>
 #include <hip_test_kernels.hh>
@@ -83,16 +82,26 @@ class MemSetKernelTest {
   }
 };
 
-static bool testhipMemsetAsyncWithKernel() {
+static bool testhipMemsetAsyncWithKernel(bool UseStrmPerThrd) {
   MemSetKernelTest<char> obj;
   constexpr char memsetval = 0x42;
 
   obj.memAllocate(memsetval);
-  for (int k = 0 ; k < ITER ; k++) {
-    hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks),
-                    dim3(threadsPerBlock), 0, obj.stream, obj.B_d, obj.C_d, N);
-    HIP_CHECK(hipMemsetAsync(obj.C_d , obj.memSetVal , N , obj.stream));
-    HIP_CHECK(hipStreamSynchronize(obj.stream));
+  for (int k = 0 ; k < ITER ; ++k) {
+    if (UseStrmPerThrd) {  // will use hipStreamPerThread stream object
+      hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks),
+                         dim3(threadsPerBlock), 0, hipStreamPerThread, obj.B_d,
+                         obj.C_d, N);
+      HIP_CHECK(hipStreamSynchronize(hipStreamPerThread));
+      HIP_CHECK(hipMemsetAsync(obj.C_d , obj.memSetVal, N, hipStreamPerThread));
+      HIP_CHECK(hipStreamSynchronize(hipStreamPerThread));
+    } else {
+      hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks),
+                      dim3(threadsPerBlock), 0, obj.stream, obj.B_d, obj.C_d,
+                      N);
+      HIP_CHECK(hipMemsetAsync(obj.C_d , obj.memSetVal , N , obj.stream));
+      HIP_CHECK(hipStreamSynchronize(obj.stream));
+    }
     HIP_CHECK(hipMemcpy(obj.A_h, obj.C_d, obj.Nbytes, hipMemcpyDeviceToHost));
 
     obj.validateExecutionOrder();
@@ -109,7 +118,7 @@ static bool testhipMemsetD32AsyncWithKernel() {
     hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks),
                     dim3(threadsPerBlock), 0, obj.stream, obj.B_d, obj.C_d, N);
     HIP_CHECK(hipMemsetD32Async((hipDeviceptr_t)obj.C_d , obj.memSetVal,
-                                                              N, obj.stream));
+                                N, obj.stream));
     HIP_CHECK(hipStreamSynchronize(obj.stream));
     HIP_CHECK(hipMemcpy(obj.A_h, obj.C_d, obj.Nbytes, hipMemcpyDeviceToHost));
 
@@ -161,7 +170,7 @@ static bool testhipMemsetD8AsyncWithKernel() {
  */
 TEST_CASE("Unit_hipMemsetAsync_VerifyExecutionWithKernel") {
   int numDevices = 0;
-  bool ret;
+  bool ret, UseStrmPerThrd = false;
 
   blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
 
@@ -172,7 +181,13 @@ TEST_CASE("Unit_hipMemsetAsync_VerifyExecutionWithKernel") {
   HIP_CHECK(hipSetDevice(devNum));
 
   SECTION("hipMemsetAsync With Kernel") {
-    ret = testhipMemsetAsyncWithKernel();
+    UseStrmPerThrd = false;
+    ret = testhipMemsetAsyncWithKernel(UseStrmPerThrd);
+    REQUIRE(ret == true);
+  }
+  SECTION("hipMemsetAsync With Kernel using hipStreamPerThread stream obj") {
+    UseStrmPerThrd = true;
+    ret = testhipMemsetAsyncWithKernel(UseStrmPerThrd);
     REQUIRE(ret == true);
   }
 

tests/src/runtimeApi/cooperativeGrps/coalesced_groups_shfl_down.cpp

@@ -0,0 +1,276 @@
+/*
+Copyright (c) 2020 - 2021 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.
+*/
+// Test Description:
+/* This test implements sum reduction kernel, first with each threads own rank
+   as input and comparing the sum with expected sum output derieved from n(n-1)/2
+   formula.
+   This sample tests functionality of intrinsics provided by thread_block_tile type,
+   shfl_down and shfl_xor.
+*/
+
+#include "test_common.h"
+#include <hip/hip_cooperative_groups.h>
+#include <stdio.h>
+#include <vector>
+
+using namespace cooperative_groups;
+
+#define ASSERT_EQUAL(lhs, rhs) assert(lhs == rhs)
+#define WAVE_SIZE 32
+
+__device__ int reduction_kernel_shfl_down(coalesced_group const& g, int val) {
+  int sz = g.size();
+
+  for (int i = sz / 2; i > 0; i >>= 1) {
+    val += g.shfl_down(val, i);
+  }
+
+  // Choose the 0'th indexed thread that holds the reduction value to return
+  if (g.thread_rank() == 0) {
+    return val;
+  }
+  // Rest of the threads return no useful values
+  else {
+    return -1;
+  }
+}
+
+__global__ void kernel_shfl_down (int * dPtr, int *dResults, int lane_delta, int cg_sizes) {
+  int id = threadIdx.x + blockIdx.x * blockDim.x;
+
+  if (id % cg_sizes == 0) {
+    coalesced_group const& g = coalesced_threads();
+    int rank = g.thread_rank();
+    int val = dPtr[rank];
+    dResults[rank] = g.shfl_down(val, lane_delta);
+    return;
+  }
+}
+
+__global__ void kernel_cg_group_partition(int* result, unsigned int tileSz, int cg_sizes) {
+
+  int id = threadIdx.x + blockIdx.x * blockDim.x;
+  if (id % cg_sizes == 0) {
+    coalesced_group  threadBlockCGTy = coalesced_threads();
+    int input, outputSum, expectedSum;
+
+    // Choose a leader thread to print the results
+    if (threadBlockCGTy.thread_rank() == 0) {
+      printf(" Creating %d groups, of tile size %d threads:\n\n",
+             (int)threadBlockCGTy.size() / tileSz, tileSz);
+    }
+
+    threadBlockCGTy.sync();
+
+    coalesced_group tiledPartition = tiled_partition(threadBlockCGTy, tileSz);
+    int threadRank = tiledPartition.thread_rank();
+
+    input = tiledPartition.thread_rank();
+
+    // (n-1)(n)/2
+    expectedSum = ((tileSz - 1) * tileSz / 2);
+
+    outputSum = reduction_kernel_shfl_down(tiledPartition, input);
+
+    if (tiledPartition.thread_rank() == 0) {
+      printf(
+          "   Sum of all ranks 0..%d in this tiledPartition group using shfl_down is %d (expected "
+          "%d)\n",
+          tiledPartition.size() - 1, outputSum, expectedSum);
+      result[threadBlockCGTy.thread_rank() / (tileSz)] = outputSum;
+    }
+    return;
+  }
+}
+
+void verifyResults(int* ptr, int expectedResult, int numTiles) {
+  for (int i = 0; i < numTiles; i++) {
+    if (ptr[i] != expectedResult) {
+      printf(" Results do not match! ");
+    }
+  }
+}
+
+void compareResults(int* cpu, int* gpu, int size) {
+  for (unsigned int i = 0; i < size / sizeof(int); i++) {
+    if (cpu[i] != gpu[i]) {
+      printf(" results do not match.");
+    }
+  }
+}
+
+void printResults(int* ptr, int size) {
+  for (int i = 0; i < size; i++) {
+    std::cout << ptr[i] << " ";
+  }
+  std::cout << '\n';
+}
+
+static void test_group_partition(unsigned int tileSz) {
+  hipError_t err;
+  int blockSize = 1;
+  int threadsPerBlock = 32;
+
+  std::vector<unsigned int> cg_sizes = {1, 2, 3};
+  for (auto i : cg_sizes) {
+
+    int numTiles = ((blockSize * threadsPerBlock) / i) / tileSz;
+    int expectedSum = ((tileSz - 1) * tileSz / 2);
+    int* expectedResult = new int[numTiles];
+
+    // numTiles = 0 when partitioning is possible. The below statement is to avoid
+    // out-of-bounds error and still evaluate failure case.
+    numTiles = (numTiles == 0) ? 1 : numTiles;
+
+    for (int i = 0; i < numTiles; i++) {
+      expectedResult[i] = expectedSum;
+    }
+
+    int* dResult = NULL;
+    int* hResult = NULL;
+
+    hipHostMalloc(&hResult, numTiles * sizeof(int), hipHostMallocDefault);
+    memset(hResult, 0, numTiles * sizeof(int));
+
+    hipMalloc(&dResult, numTiles * sizeof(int));
+
+
+    // Launch Kernel
+    hipLaunchKernelGGL(kernel_cg_group_partition, blockSize, threadsPerBlock,
+                       threadsPerBlock * sizeof(int), 0, dResult, tileSz, i);
+    err = hipDeviceSynchronize();
+    if (err != hipSuccess) {
+      fprintf(stderr, "Failed to launch kernel (error code %s)!\n", hipGetErrorString(err));
+    }
+
+
+    hipMemcpy(hResult, dResult, sizeof(int) * numTiles, hipMemcpyDeviceToHost);
+
+    verifyResults(hResult, expectedSum, numTiles);
+
+    // Free all allocated memory on host and device
+    hipFree(dResult);
+    hipFree(hResult);
+    delete[] expectedResult;
+
+    printf("\n...PASSED.\n\n");
+  }
+}
+
+static void test_shfl_down() {
+
+  std::vector<unsigned int> cg_sizes = {1, 2, 3};
+  for (auto i : cg_sizes) {
+
+    hipError_t err;
+    int blockSize = 1;
+    int threadsPerBlock = WAVE_SIZE;
+
+    int totalThreads = blockSize * threadsPerBlock;
+    int group_size = totalThreads / i;
+    int group_size_in_bytes = group_size * sizeof(int);
+
+    int* hPtr = NULL;
+    int* dPtr = NULL;
+    int* dResults = NULL;
+    int lane_delta = rand() % group_size;
+    std::cout << "Testing coalesced_groups shfl_down with lane_delta " << lane_delta << "and group size "
+              << WAVE_SIZE << '\n' << std::endl;
+
+    int arrSize = blockSize * threadsPerBlock * sizeof(int);
+
+    hipHostMalloc(&hPtr, arrSize);
+    // Fill up the array
+    for (int i = 0; i < WAVE_SIZE; i++) {
+      hPtr[i] = rand() % 1000;
+    }
+
+    int* cpuResultsArr = (int*)malloc(group_size_in_bytes);
+    for (int i = 0; i < group_size; i++) {
+      cpuResultsArr[i] = (i + lane_delta >= group_size) ? hPtr[i] : hPtr[i + lane_delta];
+    }
+    //printf("Array passed to GPU for computation\n");
+    //printResults(hPtr, WAVE_SIZE);
+    hipMalloc(&dPtr, group_size_in_bytes);
+    hipMalloc(&dResults, group_size_in_bytes);
+
+    hipMemcpy(dPtr, hPtr, group_size_in_bytes, hipMemcpyHostToDevice);
+    // Launch Kernel
+    hipLaunchKernelGGL(kernel_shfl_down, blockSize, threadsPerBlock,
+                       threadsPerBlock * sizeof(int), 0, dPtr, dResults, lane_delta, i);
+    hipMemcpy(hPtr, dResults, group_size_in_bytes, hipMemcpyDeviceToHost);
+    err = hipDeviceSynchronize();
+    if (err != hipSuccess) {
+      fprintf(stderr, "Failed to launch kernel (error code %s)!\n", hipGetErrorString(err));
+    }
+    //printf("GPU results: \n");
+    //printResults(hPtr, WAVE_SIZE);
+    //printf("Printing cpu to be verified array\n");
+    //printResults(cpuResultsArr, WAVE_SIZE);
+
+    compareResults(hPtr, cpuResultsArr, group_size_in_bytes);
+    std::cout << "Results verified!\n";
+
+    hipFree(hPtr);
+    hipFree(dPtr);
+    free(cpuResultsArr);
+  }
+}
+
+
+
+int main() {
+  // Use default device for validating the test
+  int deviceId;
+  ASSERT_EQUAL(hipGetDevice(&deviceId), hipSuccess);
+  hipDeviceProp_t deviceProperties;
+  ASSERT_EQUAL(hipGetDeviceProperties(&deviceProperties, deviceId), hipSuccess);
+  int maxThreadsPerBlock = deviceProperties.maxThreadsPerBlock;
+
+  if (!deviceProperties.cooperativeLaunch) {
+    std::cout << "info: Device doesn't support cooperative launch! skipping the test!\n";
+    if (hip_skip_tests_enabled()) {
+      return hip_skip_retcode();
+    } else {
+      passed();
+    }
+    return 0;
+  }
+
+  // Test shfl_down with random group sizes
+    for (int i = 0; i < 100; i++) {
+      test_shfl_down();
+    }
+
+  std::cout << "Testing static tiled_partition for different tile sizes using shfl_down"
+            << std::endl;
+
+  int testNo = 1;
+  std::vector<unsigned int> tileSizes = {2, 4, 8, 16, 32};
+  for (auto i : tileSizes) {
+    std::cout << "TEST " << testNo << ":" << '\n' << std::endl;
+    test_group_partition(i);
+    testNo++;
+  }
+
+  passed();
+}

tests/src/runtimeApi/cooperativeGrps/coalesced_groups_shfl_up.cpp

@@ -0,0 +1,260 @@
+/*
+Copyright (c) 2020 - 2021 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.
+*/
+// Test Description:
+/* This test implements prefix sum(scan) kernel, first with each threads own rank
+   as input and comparing the sum with expected serial summation output on CPU.
+
+   This sample tests functionality of intrinsics provided by coalesced_group,
+   shfl_up.
+*/
+#include "test_common.h"
+#include <hip/hip_cooperative_groups.h>
+#include <stdio.h>
+#include <vector>
+
+using namespace cooperative_groups;
+
+#define ASSERT_EQUAL(lhs, rhs) assert(lhs == rhs)
+#define WAVE_SIZE 32
+__device__ int prefix_sum_kernel(coalesced_group const& g, int val) {
+  int sz = g.size();
+  for (int i = 1; i < sz; i <<= 1) {
+    int temp = g.shfl_up(val, i);
+
+    if (g.thread_rank() >= i) {
+      val += temp;
+    }
+  }
+  return val;
+}
+
+__global__ void kernel_shfl_up (int * dPtr, int *dResults, int lane_delta, int cg_sizes) {
+  int id = threadIdx.x + blockIdx.x * blockDim.x;
+
+  if (id % cg_sizes == 0) {
+    coalesced_group g = coalesced_threads();
+    int rank = g.thread_rank();
+    int val = dPtr[rank];
+    dResults[rank] = g.shfl_up(val, lane_delta);
+  return;
+  }
+
+}
+
+__global__ void kernel_cg_group_partition(int* dPtr, unsigned int tileSz, int cg_sizes) {
+
+  int id = threadIdx.x + blockIdx.x * blockDim.x;
+  if (id % cg_sizes == 0) {
+    coalesced_group threadBlockCGTy = coalesced_threads();
+    int input, outputSum;
+
+    // we pass its own thread rank as inputs
+    input = threadBlockCGTy.thread_rank();
+
+    // Choose a leader thread to print the results
+    if (threadBlockCGTy.thread_rank() == 0) {
+      printf(" Creating %d groups, of tile size %d threads:\n\n",
+             (int)threadBlockCGTy.size() / tileSz, tileSz);
+    }
+
+    threadBlockCGTy.sync();
+
+    coalesced_group tiledPartition = tiled_partition(threadBlockCGTy, tileSz);
+
+    input = tiledPartition.thread_rank();
+
+    outputSum = prefix_sum_kernel(tiledPartition, input);
+
+    // Update the result array with the corresponsing prefix sum
+    dPtr[threadBlockCGTy.thread_rank()] = outputSum;
+    return;
+  }
+}
+
+void serialScan(int* ptr, int size) {
+  // Fill up the array
+  for (int i = 0; i < size; i++) {
+    ptr[i] = i;
+  }
+
+  int acc = 0;
+  for (int i = 0; i < size; i++) {
+    acc = acc + ptr[i];
+    ptr[i] = acc;
+  }
+}
+
+void printResults(int* ptr, int size) {
+  for (int i = 0; i < size; i++) {
+    std::cout << ptr[i] << " ";
+  }
+  std::cout << '\n';
+}
+
+void verifyResults(int* cpu, int* gpu, int size) {
+  for (unsigned int i = 0; i < size / sizeof(int); i++) {
+    if (cpu[i] != gpu[i]) {
+      printf(" Results do not match.");
+    }
+  }
+}
+
+static void test_group_partition(unsigned tileSz) {
+  hipError_t err;
+  int blockSize = 1;
+  int threadsPerBlock = WAVE_SIZE;
+
+  int* hPtr = NULL;
+  int* dPtr = NULL;
+  int* cpuPrefixSum = NULL;
+
+  std::vector<unsigned int> cg_sizes = {1, 2, 3};
+  for (auto i : cg_sizes) {
+
+    int arrSize = blockSize * threadsPerBlock * sizeof(int);
+
+    hipHostMalloc(&hPtr, arrSize);
+    hipMalloc(&dPtr, arrSize);
+
+    // Launch Kernel
+    hipLaunchKernelGGL(kernel_cg_group_partition, blockSize, threadsPerBlock,
+                     threadsPerBlock * sizeof(int), 0, dPtr, tileSz, i);
+    hipMemcpy(hPtr, dPtr, arrSize, hipMemcpyDeviceToHost);
+    err = hipDeviceSynchronize();
+    if (err != hipSuccess) {
+      fprintf(stderr, "Failed to launch kernel (error code %s)!\n", hipGetErrorString(err));
+    }
+
+    cpuPrefixSum = new int[tileSz];
+    serialScan(cpuPrefixSum, tileSz);
+    //std::cout << "\nPrefix sum results on CPU\n";
+    //printResults(cpuPrefixSum, tileSz);
+
+    //std::cout << "\nPrefix sum results on GPU\n";
+    //printResults(hPtr, tileSz);
+    std::cout << "\n";
+    verifyResults(hPtr, cpuPrefixSum, tileSz);
+    std::cout << "Results verified!\n";
+
+    delete[] cpuPrefixSum;
+    hipFree(hPtr);
+    hipFree(dPtr);
+  }
+}
+
+static void test_shfl_up() {
+
+  std::vector<unsigned int> cg_sizes = {1, 2, 3};
+  for (auto i : cg_sizes) {
+
+    hipError_t err;
+    int blockSize = 1;
+
+    int threadsPerBlock = WAVE_SIZE;
+    int totalThreads = blockSize * threadsPerBlock;
+    int group_size = totalThreads / i;
+    int group_size_in_bytes = group_size * sizeof(int);
+
+    int* hPtr = NULL;
+    int* dPtr = NULL;
+    int* dResults = NULL;
+    int lane_delta =  (rand() % group_size);
+
+    std::cout << "Testing coalesced_groups shfl_up with lane_delta " << lane_delta
+              << " and group size " << WAVE_SIZE << '\n' << std::endl;
+
+    int arrSize = blockSize * threadsPerBlock * sizeof(int);
+
+    hipHostMalloc(&hPtr, arrSize);
+    // Fill up the array
+    for (int i = 0; i < WAVE_SIZE; i++) {
+      hPtr[i] = rand() % 1000;
+    }
+    //printResults(hPtr, WAVE_SIZE);
+
+    int* cpuResultsArr = (int*)malloc(group_size_in_bytes);
+    for (int i = 0; i < group_size; i++) {
+      cpuResultsArr[i] = (i <= (lane_delta - 1)) ?  hPtr[i] : hPtr[i - lane_delta];
+    }
+
+    //printf("Printing cpu results arr\n");
+    //printResults(cpuResultsArr, WAVE_SIZE);
+
+    hipMalloc(&dPtr, group_size_in_bytes);
+    hipMalloc(&dResults, group_size_in_bytes);
+
+    hipMemcpy(dPtr, hPtr, group_size_in_bytes, hipMemcpyHostToDevice);
+    // Launch Kernel
+    hipLaunchKernelGGL(kernel_shfl_up, blockSize, threadsPerBlock,
+                       threadsPerBlock * sizeof(int), 0, dPtr, dResults, lane_delta, i);
+    hipMemcpy(hPtr, dResults, group_size_in_bytes, hipMemcpyDeviceToHost);
+    err = hipDeviceSynchronize();
+    if (err != hipSuccess) {
+      fprintf(stderr, "Failed to launch kernel (error code %s)!\n", hipGetErrorString(err));
+    }
+    //printf("GPU computation array :\n");
+    //printResults(hPtr, WAVE_SIZE);
+
+    verifyResults(hPtr, cpuResultsArr, group_size_in_bytes);
+    std::cout << "Results verified!\n";
+
+    hipFree(hPtr);
+    hipFree(dPtr);
+    free(cpuResultsArr);
+  }
+}
+
+int main() {
+  // Use default device for validating the test
+  int deviceId;
+  ASSERT_EQUAL(hipGetDevice(&deviceId), hipSuccess);
+  hipDeviceProp_t deviceProperties;
+  ASSERT_EQUAL(hipGetDeviceProperties(&deviceProperties, deviceId), hipSuccess);
+  int maxThreadsPerBlock = deviceProperties.maxThreadsPerBlock;
+
+  if (!deviceProperties.cooperativeLaunch) {
+    std::cout << "info: Device doesn't support cooperative launch! skipping the test!\n";
+    if (hip_skip_tests_enabled()) {
+      return hip_skip_retcode();
+    } else {
+      passed();
+    }
+    return 0;
+  }
+
+    for (int i = 0; i < 100; i++) {
+      test_shfl_up();
+    }
+
+  std::cout << "Testing coalesced_groups partitioning and shfl_up" << '\n' << std::endl;
+
+  int testNo = 1;
+  std::vector<unsigned int> tileSizes = {2, 4, 8, 16, 32};
+  for (auto i : tileSizes) {
+    std::cout << "TEST " << testNo << ":" << '\n' << std::endl;
+    test_group_partition(i);
+    testNo++;
+  }
+  passed();
+}
+
+/* Kogge-Stone algorithm */

tests/src/runtimeApi/cooperativeGrps/simple_coalesced_groups.cpp

@@ -0,0 +1,583 @@
+/*
+Copyright (c) 2020 - 2021 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.
+*/
+
+// Test Description:
+/* This test implements sum reduction kernel, first with each threads own rank
+   as input and comparing the sum with expected sum output derieved from n(n-1)/2
+   formula. The second part, partitions this parent group into child subgroups
+   a.k.a tiles using using tiled_partition() collective operation. This can be called
+   with a static tile size, passed in templated non-type variable-tiled_partition<tileSz>,
+   or in runtime as tiled_partition(thread_group parent, tileSz). This test covers both these
+   cases.
+   This test tests functionality of cg group partitioning, (static and dynamic) and its respective
+   API's size(), thread_rank(), and sync().
+*/
+
+#include "test_common.h"
+#include <hip/hip_cooperative_groups.h>
+#include <stdio.h>
+#include <vector>
+
+using namespace cooperative_groups;
+
+#define ASSERT_EQUAL(lhs, rhs) assert(lhs == rhs)
+
+#define NUM_ELEMS 10000000
+#define NUM_THREADS_PER_BLOCK 512
+#define WAVE_SIZE 32
+
+/* Test coalesced group's functionality.
+ *
+ */
+
+__device__ int atomicAggInc(int *ptr) {
+   coalesced_group g = coalesced_threads();
+   int prev;
+   // elect the first active thread to perform atomic add
+   if (g.thread_rank() == 0) {
+     prev = atomicAdd(ptr, g.size());
+   }
+   // broadcast previous value within the warp
+   // and add each active thread’s rank to it
+   prev = g.thread_rank() + g.shfl(prev, 0);
+   return prev;
+}
+
+__global__ void kernel_shfl (int * dPtr, int *dResults, int srcLane, int cg_sizes) {
+  int id = threadIdx.x + blockIdx.x * blockDim.x;
+  if (id % cg_sizes == 0) {
+    coalesced_group const& g = coalesced_threads();
+    int rank = g.thread_rank();
+    int val = dPtr[rank];
+    dResults[rank] = g.shfl(val, srcLane);
+    return;
+  }
+}
+
+__global__ void kernel_shfl_any_to_any (int *randVal, int *dsrcArr, int *dResults, int cg_sizes) {
+
+ int id = threadIdx.x + blockIdx.x * blockDim.x;
+
+ if (id % cg_sizes == 0) {
+    coalesced_group const& g = coalesced_threads();
+    int rank = g.thread_rank();
+    int val = randVal[rank];
+    dResults[rank] = g.shfl(val, dsrcArr[rank]);
+    return;
+  }
+
+}
+
+__global__ void filter_arr(int *dst, int *nres, const int *src, int n) {
+  int id = threadIdx.x + blockIdx.x * blockDim.x;
+
+  for (int i = id; i < n; i += gridDim.x * blockDim.x) {
+    if (src[i] > 0) dst[atomicAggInc(nres)] = src[i];
+  }
+}
+
+/* Parallel reduce kernel.
+ *
+ * Step complexity: O(log n)
+ * Work complexity: O(n)
+ *
+ * Note: This kernel works only with power of 2 input arrays.
+ */
+__device__ int reduction_kernel(coalesced_group g, int* x, int val) {
+  int lane = g.thread_rank();
+  int sz = g.size();
+
+  for (int i = g.size() / 2; i > 0; i /= 2) {
+    // use lds to store the temporary result
+    x[lane] = val;
+    // Ensure all the stores are completed.
+    g.sync();
+
+    if (lane < i) {
+      val += x[lane + i];
+    }
+    // It must work on one tiled thread group at a time,
+    // and it must make sure all memory operations are
+    // completed before moving to the next stride.
+    // sync() here just does that.
+    g.sync();
+  }
+
+  // Choose the 0'th indexed thread that holds the reduction value to return
+  if (g.thread_rank() == 0) {
+    return val;
+  }
+  // Rest of the threads return no useful values
+  else {
+    return -1;
+  }
+}
+
+__global__ void kernel_cg_coalesced_group_partition(unsigned int tileSz, int* result,
+                                                  bool isGlobalMem, int* globalMem, int cg_sizes) {
+
+  int id = threadIdx.x + blockIdx.x * blockDim.x;
+  if (id % cg_sizes == 0) {
+    coalesced_group threadBlockCGTy = coalesced_threads();
+    int threadBlockGroupSize = threadBlockCGTy.size();
+
+    int* workspace = NULL;
+
+    if (isGlobalMem) {
+      workspace = globalMem;
+    } else {
+      // Declare a shared memory
+      extern __shared__ int sharedMem[];
+      workspace = sharedMem;
+    }
+
+    int input, outputSum, expectedOutput;
+
+    // input to reduction, for each thread, is its' rank in the group
+    input = threadBlockCGTy.thread_rank();
+
+    expectedOutput = (threadBlockGroupSize - 1) * threadBlockGroupSize / 2;
+
+    outputSum = reduction_kernel(threadBlockCGTy, workspace, input);
+
+    if (threadBlockCGTy.thread_rank() == 0) {
+      printf(" Sum of all ranks 0..%d in coalesced_group is %d\n\n",
+             (int)threadBlockCGTy.size() - 1, outputSum);
+      printf(" Creating %d groups, of tile size %d threads:\n\n",
+             (int)threadBlockCGTy.size() / tileSz, tileSz);
+    }
+
+    threadBlockCGTy.sync();
+
+    coalesced_group tiledPartition = tiled_partition(threadBlockCGTy, tileSz);
+
+    // This offset allows each group to have its own unique area in the workspace array
+    int workspaceOffset = threadBlockCGTy.thread_rank() - tiledPartition.thread_rank();
+
+    outputSum = reduction_kernel(tiledPartition, workspace + workspaceOffset, input);
+
+    if (tiledPartition.thread_rank() == 0) {
+      printf(
+          "   Sum of all ranks 0..%d in this tiledPartition group is %d. Corresponding parent thread "
+          "rank: %d\n",
+          tiledPartition.size() - 1, outputSum, input);
+
+        result[input / (tileSz)] = outputSum;
+    }
+    return;
+  }
+}
+
+__global__ void kernel_coalesced_active_groups() {
+  thread_block threadBlockCGTy = this_thread_block();
+  int threadBlockGroupSize = threadBlockCGTy.size();
+
+  // input to reduction, for each thread, is its' rank in the group
+  int input = threadBlockCGTy.thread_rank();
+
+  if (threadBlockCGTy.thread_rank() == 0) {
+    printf(" Creating odd and even set of active thread groups based on branch divergence\n\n");
+  }
+
+  threadBlockCGTy.sync();
+
+  // Group all active odd threads
+  if (threadBlockCGTy.thread_rank() % 2) {
+    coalesced_group activeOdd = coalesced_threads();
+
+    if (activeOdd.thread_rank() == 0) {
+      printf(" ODD: Size of odd set of active threads is %d."
+             " Corresponding parent thread_rank is %d.\n\n",
+               activeOdd.size(), threadBlockCGTy.thread_rank());
+    }
+  }
+  else { // Group all active even threads
+    coalesced_group activeEven = coalesced_threads();
+
+    if (activeEven.thread_rank() == 0) {
+      printf(" EVEN: Size of even set of active threads is %d."
+             " Corresponding parent thread_rank is %d.",
+               activeEven.size(), threadBlockCGTy.thread_rank());
+    }
+  }
+  return;
+}
+
+void printResults(int* ptr, int size) {
+  for (int i = 0; i < size; i++) {
+    std::cout << ptr[i] << " ";
+  }
+  std::cout << '\n';
+}
+
+void compareResults(int* cpu, int* gpu, int size) {
+  for (unsigned int i = 0; i < size / sizeof(int); i++) {
+    if (cpu[i] != gpu[i]) {
+      printf(" results do not match.");
+    }
+  }
+}
+
+static void test_active_threads_grouping() {
+  hipError_t err;
+  int blockSize = 1;
+  int threadsPerBlock = WAVE_SIZE;
+
+  // Launch Kernel
+    hipLaunchKernelGGL(kernel_coalesced_active_groups, blockSize, threadsPerBlock, 0, 0);
+
+    err = hipDeviceSynchronize();
+    if (err != hipSuccess) {
+      fprintf(stderr, "Failed to launch kernel (error code %s)!\n", hipGetErrorString(err));
+    }
+  printf("\n...PASSED.\n\n");
+}
+
+// Search if the sum exists in the expected results array
+void verifyResults(int* hPtr, int* dPtr, int size) {
+  int i = 0, j = 0;
+  for (i = 0; i < size; i++) {
+    for (j = 0; j < size; j++) {
+      if (hPtr[i] == dPtr[j]) {
+        break;
+      }
+    }
+    if (j == size) {
+      printf(" Result verification failed!");
+    }
+  }
+}
+
+
+static void test_group_partition(unsigned int tileSz, bool useGlobalMem) {
+  hipError_t err;
+  int blockSize = 1;
+  int threadsPerBlock = WAVE_SIZE;
+
+  std::vector<unsigned int> cg_sizes = {1, 2, 3};
+  for (auto i : cg_sizes) {
+
+    int numTiles = ((blockSize * threadsPerBlock) / i) / tileSz;
+
+    // numTiles = 0 when partitioning is possible. The below statement is to avoid
+    // out-of-bounds error and still evaluate failure case.
+    numTiles = (numTiles == 0) ? 1 : numTiles;
+
+    // Build an array of expected reduction sum output on the host
+    // based on the sum of their respective thread ranks to use for verification
+    int* expectedSum = new int[numTiles];
+    int temp = 0, sum = 0;
+    for (int i = 1; i <= numTiles; i++) {
+      sum = temp;
+      temp = (((tileSz * i) - 1) * (tileSz * i)) / 2;
+      expectedSum[i-1] = temp - sum;
+    }
+
+    int* dResult = NULL;
+    hipMalloc(&dResult, sizeof(int) * numTiles);
+
+    int* globalMem = NULL;
+    if (useGlobalMem) {
+      hipMalloc((void**)&globalMem, threadsPerBlock * sizeof(int));
+    }
+
+    int* hResult = NULL;
+    hipHostMalloc(&hResult, numTiles * sizeof(int), hipHostMallocDefault);
+    memset(hResult, 0, numTiles * sizeof(int));
+
+    // Launch Kernel
+    if (useGlobalMem) {
+      hipLaunchKernelGGL(kernel_cg_coalesced_group_partition, blockSize, threadsPerBlock, 0, 0, tileSz,
+                         dResult, useGlobalMem, globalMem, i);
+
+      err = hipDeviceSynchronize();
+      if (err != hipSuccess) {
+        fprintf(stderr, "Failed to launch kernel (error code %s)!\n", hipGetErrorString(err));
+      }
+    } else {
+      hipLaunchKernelGGL(kernel_cg_coalesced_group_partition, blockSize, threadsPerBlock,
+                         threadsPerBlock * sizeof(int), 0, tileSz, dResult, useGlobalMem, globalMem, i);
+
+      err = hipDeviceSynchronize();
+      if (err != hipSuccess) {
+        fprintf(stderr, "Failed to launch kernel (error code %s)!\n", hipGetErrorString(err));
+      }
+    }
+
+    hipMemcpy(hResult, dResult, numTiles * sizeof(int), hipMemcpyDeviceToHost);
+    verifyResults(expectedSum, hResult, numTiles);
+    // Free all allocated memory on host and device
+    hipFree(dResult);
+    hipFree(hResult);
+    if (useGlobalMem) {
+      hipFree(globalMem);
+    }
+    delete[] expectedSum;
+
+    printf("\n...PASSED.\n\n");
+  }
+}
+static void test_shfl_any_to_any() {
+
+  std::vector<unsigned int> cg_sizes = {1, 2, 3};
+  for (auto i : cg_sizes) {
+
+    hipError_t err;
+    int blockSize = 1;
+    int threadsPerBlock = WAVE_SIZE;
+
+    int totalThreads = blockSize * threadsPerBlock;
+    int group_size = (totalThreads + i - 1) / i;
+    int group_size_in_bytes = group_size * sizeof(int);
+
+    int* hPtr = NULL;
+    int* dPtr = NULL;
+    int* dsrcArr = NULL;
+    int* dResults = NULL;
+    int* srcArr = (int*)malloc(group_size_in_bytes);
+    int* srcArrCpu = (int*)malloc(group_size_in_bytes);
+
+    std::cout << "Testing coalesced_groups shfl any-to-any\n" <<std::endl;
+
+    int arrSize = blockSize * threadsPerBlock * sizeof(int);
+
+    hipHostMalloc(&hPtr, arrSize);
+    // Fill up the array
+    for (int i = 0; i < WAVE_SIZE; i++) {
+      hPtr[i] = rand() % 1000;
+    }
+
+    // Fill up the random array
+    for (int i = 0; i < group_size; i++) {
+      srcArr[i] = rand() % 1000;
+      srcArrCpu[i] = srcArr[i] % group_size;
+    }
+
+    /* Fill cpu results array so that we can verify with gpu computation */
+    int* cpuResultsArr = (int*)malloc(group_size_in_bytes);
+    for(int i = 0; i < group_size; i++) {
+      cpuResultsArr[i] = hPtr[srcArrCpu[i]];
+    }
+
+    //printf("Array passed to GPU for computation\n");
+    //printResults(hPtr, WAVE_SIZE);
+    hipMalloc(&dPtr, group_size_in_bytes);
+    hipMalloc(&dResults, group_size_in_bytes);
+
+    hipMalloc(&dsrcArr, group_size_in_bytes);
+    hipMemcpy(dsrcArr, srcArr, group_size_in_bytes, hipMemcpyHostToDevice);
+
+    hipMemcpy(dPtr, hPtr, group_size_in_bytes, hipMemcpyHostToDevice);
+    // Launch Kernel
+    hipLaunchKernelGGL(kernel_shfl_any_to_any, blockSize, threadsPerBlock,
+                       threadsPerBlock * sizeof(int), 0 , dPtr, dsrcArr, dResults, i);
+    hipMemcpy(hPtr, dResults, group_size_in_bytes, hipMemcpyDeviceToHost);
+    err = hipDeviceSynchronize();
+    if (err != hipSuccess) {
+      fprintf(stderr, "Failed to launch kernel (error code %s)!\n", hipGetErrorString(err));
+    }
+    //printf("GPU results: \n");
+    //printResults(hPtr, group_size);
+    //printf("Printing cpu to be verified array\n");
+    //printResults(cpuResultsArr, group_size);
+    //printf("Printing srcLane array that was passed\n");
+    //printResults(srcArr, group_size);
+    //printf("Printing srcLane array on the CPU\n");
+    //printResults(srcArrCpu, group_size);
+    compareResults(hPtr, cpuResultsArr, group_size_in_bytes);
+    std::cout << "Results verified!\n";
+
+    hipFree(hPtr);
+    hipFree(dPtr);
+    free(srcArr);
+    free(srcArrCpu);
+    free(cpuResultsArr);
+  }
+}
+static void test_shfl_broadcast() {
+
+  std::vector<unsigned int> cg_sizes = {1, 2, 3};
+  for (auto i : cg_sizes) {
+
+    hipError_t err;
+    int blockSize = 1;
+    int threadsPerBlock = WAVE_SIZE;
+
+    int totalThreads = blockSize * threadsPerBlock;
+    int group_size = (totalThreads + i - 1) / i;
+    int group_size_in_bytes = group_size * sizeof(int);
+
+    int* hPtr = NULL;
+    int* dPtr = NULL;
+    int* dResults = NULL;
+    int srcLane = rand() % 1000;
+    int srcLaneCpu = 0;
+    std::cout << "Testing coalesced_groups shfl with srcLane " << srcLane << '\n'
+              << " and group size " << i <<std::endl;
+
+    int arrSize = blockSize * threadsPerBlock * sizeof(int);
+
+    hipHostMalloc(&hPtr, arrSize);
+    // Fill up the array
+    for (int i = 0; i < WAVE_SIZE; i++) {
+      hPtr[i] = rand() % 1000;
+    }
+
+
+    /* Fill cpu results array so that we can verify with gpu computation */
+    srcLaneCpu = hPtr[srcLane % group_size];
+
+    int* cpuResultsArr = (int*)malloc(sizeof(int) * group_size);
+    for (int i = 0; i < group_size; i++) {
+      cpuResultsArr[i] = srcLaneCpu;
+    }
+    printf("Array passed to GPU for computation\n");
+    printResults(hPtr, WAVE_SIZE);
+    hipMalloc(&dPtr, group_size_in_bytes);
+    hipMalloc(&dResults, group_size_in_bytes);
+
+    hipMemcpy(dPtr, hPtr, group_size_in_bytes, hipMemcpyHostToDevice);
+    // Launch Kernel
+    hipLaunchKernelGGL(kernel_shfl, blockSize, threadsPerBlock,
+                       threadsPerBlock * sizeof(int), 0, dPtr, dResults, srcLane, i);
+    hipMemcpy(hPtr, dResults, group_size_in_bytes, hipMemcpyDeviceToHost);
+    err = hipDeviceSynchronize();
+    if (err != hipSuccess) {
+      fprintf(stderr, "Failed to launch kernel (error code %s)!\n", hipGetErrorString(err));
+    }
+    printf("GPU results: \n");
+    printResults(hPtr, group_size);
+    printf("Printing cpu to be verified array\n");
+    printResults(cpuResultsArr, group_size);
+
+    compareResults(hPtr, cpuResultsArr, group_size_in_bytes);
+    std::cout << "Results verified!\n";
+
+    hipFree(hPtr);
+    hipFree(dPtr);
+    free(cpuResultsArr);
+  }
+}
+
+int main() {
+  // Use default device for validating the test
+  int deviceId;
+  ASSERT_EQUAL(hipGetDevice(&deviceId), hipSuccess);
+  hipDeviceProp_t deviceProperties;
+  ASSERT_EQUAL(hipGetDeviceProperties(&deviceProperties, deviceId), hipSuccess);
+  int maxThreadsPerBlock = deviceProperties.maxThreadsPerBlock;
+
+  if (!deviceProperties.cooperativeLaunch) {
+    std::cout << "info: Device doesn't support cooperative launch! skipping the test!\n";
+    if (hip_skip_tests_enabled()) {
+      return hip_skip_retcode();
+    } else {
+      passed();
+    }
+  }
+
+  std::cout << "Now testing coalesced_groups" << '\n' << std::endl;
+
+  int *data_to_filter, *filtered_data, nres = 0;
+  int *d_data_to_filter, *d_filtered_data, *d_nres;
+
+  int numOfBuckets = 5;
+
+  data_to_filter = reinterpret_cast<int *>(malloc(sizeof(int) * NUM_ELEMS));
+
+  // Generate input data.
+  for (int i = 0; i < NUM_ELEMS; i++) {
+    data_to_filter[i] = rand() % numOfBuckets;
+  }
+
+
+  hipMalloc(&d_data_to_filter, sizeof(int) * NUM_ELEMS);
+  hipMalloc(&d_filtered_data, sizeof(int) * NUM_ELEMS);
+  hipMalloc(&d_nres, sizeof(int));
+
+  hipMemcpy(d_data_to_filter, data_to_filter,
+                             sizeof(int) * NUM_ELEMS, hipMemcpyHostToDevice);
+  hipMemset(d_nres, 0, sizeof(int));
+
+  dim3 dimBlock(NUM_THREADS_PER_BLOCK, 1, 1);
+  dim3 dimGrid((NUM_ELEMS / NUM_THREADS_PER_BLOCK) + 1, 1, 1);
+
+  filter_arr<<<dimGrid, dimBlock>>>(d_filtered_data, d_nres, d_data_to_filter,
+                                    NUM_ELEMS);
+
+
+  hipMemcpy(&nres, d_nres, sizeof(int), hipMemcpyDeviceToHost);
+
+  filtered_data = reinterpret_cast<int *>(malloc(sizeof(int) * nres));
+
+  hipMemcpy(filtered_data, d_filtered_data, sizeof(int) * nres,
+                             hipMemcpyDeviceToHost);
+
+  int *host_filtered_data =
+      reinterpret_cast<int *>(malloc(sizeof(int) * NUM_ELEMS));
+
+  // Generate host output with host filtering code.
+  int host_flt_count = 0;
+  for (int i = 0; i < NUM_ELEMS; i++) {
+    if (data_to_filter[i] > 0) {
+      host_filtered_data[host_flt_count++] = data_to_filter[i];
+    }
+  }
+
+  printf("\nWarp Aggregated Atomics %s \n",
+         (host_flt_count == nres) ? "PASSED" : "FAILED");
+
+  // Now, testing shfl collective
+  std::cout << "Now testing shfl collective as a broadcast" << '\n' << std::endl;
+
+  for (int i = 0; i < 100; i++) {
+    test_shfl_broadcast();
+  }
+
+
+  // Now, testing shfl collective
+  std::cout << "Now testing shfl operations any-to-any member lanes" << '\n' << std::endl;
+
+  for (int i = 0; i < 100; i++) {
+    test_shfl_any_to_any();
+  }
+
+  // Now, pass a already coalesced_group that was partitioned
+  /* Test coalesced group partitioning */
+  std::cout << "Now testing coalesced_groups partitioning" << '\n' << std::endl;
+
+  int testNo = 1;
+  for (int memTy = 0; memTy < 2; memTy++) {
+    std::vector<unsigned int> tileSizes = {2, 4, 8, 16, 32};
+    for (auto i : tileSizes) {
+      std::cout << "TEST " << testNo << ":" << '\n' << std::endl;
+      test_group_partition(i, memTy);
+      testNo++;
+    }
+  }
+
+  std::cout << "Now grouping active threads based on branch divergence" << '\n' << std::endl;
+  test_active_threads_grouping();
+
+  passed();
+  return 0;
+}

tests/src/runtimeApi/graph/hipGraph.cpp

@@ -105,6 +105,25 @@ bool hipWithoutGraphs(float* inputVec_h, float* inputVec_d, double* outputVec_d,
   return true;
 }
 
+typedef struct callBackData {
+  const char* fn_name;
+  double* data;
+} callBackData_t;
+double result_gpu = 0.0;
+void myHostNodeCallback(void* data) {
+  static int iter = 0;
+  iter++;
+  // Check status of GPU after stream operations are done
+  callBackData_t* tmp = (callBackData_t*)(data);
+  // checkCudaErrors(tmp->status);
+  double* result = (double*)(tmp->data);
+  char* function = (char*)(tmp->fn_name);
+  if (iter == GRAPH_LAUNCH_ITERATIONS)
+    printf("[%s] Host callback final reduced sum = %lf\n", function, *result);
+  result_gpu = *result;
+  *result = 0.0;  // reset the result
+}
+
 bool hipGraphsUsingStreamCapture(float* inputVec_h, float* inputVec_d, double* outputVec_d,
                                  double* result_d, size_t inputSize, size_t numOfBlocks) {
   hipStream_t stream1, stream2, stream3, streamForGraph;
@@ -237,6 +256,16 @@ bool hipGraphsManual(float* inputVec_h, float* inputVec_d, double* outputVec_d,
   nodeDependencies.clear();
   nodeDependencies.push_back(memcpyNode);
   hipGraphNode_t hostNode;
+  hipHostNodeParams hostParams = {0};
+  hostParams.fn = myHostNodeCallback;
+  callBackData_t hostFnData;
+  hostFnData.data = &result_h;
+  hostFnData.fn_name = "hipGraphsManual";
+  hostParams.userData = &hostFnData;
+
+  HIPCHECK(hipGraphAddHostNode(&hostNode, graph, nodeDependencies.data(), nodeDependencies.size(),
+                               &hostParams));
+
   hipGraphExec_t graphExec;
   hipGraphNode_t* nodes = NULL;
   size_t numNodes = 0;
@@ -266,8 +295,8 @@ bool hipGraphsManual(float* inputVec_h, float* inputVec_d, double* outputVec_d,
   for (int i = 0; i < inputSize; i++) {
     result_h_cpu += inputVec_h[i];
   }
-  if (result_h_cpu != result_h) {
-    printf("Final reduced sum = %lf %lf\n", result_h_cpu, result_h);
+  if (result_h_cpu != result_gpu) {
+    printf("Final reduced sum = %lf %lf\n", result_h_cpu, result_gpu);
     return false;
   }
   return true;
@@ -304,4 +333,4 @@ int main(int argc, char** argv) {
     failed("Failed during hipGraph with capture\n");
   }
   passed();
-}
+}