diff --git a/catch/include/memcpy3d_tests_common.hh b/catch/include/memcpy3d_tests_common.hh
index 84d0fc517b..acb549c94a 100644
--- a/catch/include/memcpy3d_tests_common.hh
+++ b/catch/include/memcpy3d_tests_common.hh
@@ -23,7 +23,7 @@ THE SOFTWARE.
 #pragma once
 #pragma clang diagnostic ignored "-Wmissing-field-initializers"
 #pragma clang diagnostic ignored "-Wunused-lambda-capture"
-
+#pragma clang diagnostic ignored "-Wunused-parameter"
 #include <variant>
 
 #include <hip_test_common.hh>
@@ -44,8 +44,9 @@ static inline hipMemcpyKind ReverseMemcpyDirection(const hipMemcpyKind direction
   }
 };
 
-static hipMemcpy3DParms GetMemcpy3DParms(PtrVariant dst_ptr, hipPos dst_pos, PtrVariant src_ptr,
-                                         hipPos src_pos, hipExtent extent, hipMemcpyKind kind) {
+static inline hipMemcpy3DParms GetMemcpy3DParms(PtrVariant dst_ptr, hipPos dst_pos,
+                                                PtrVariant src_ptr, hipPos src_pos,
+                                                hipExtent extent, hipMemcpyKind kind) {
   hipMemcpy3DParms parms = {0};
   if (std::holds_alternative<hipArray_t>(dst_ptr)) {
     parms.dstArray = std::get<hipArray_t>(dst_ptr);
@@ -185,7 +186,7 @@ void Memcpy3DDeviceToDeviceShell(F memcpy_func, hipStream_t kernel_stream = null
     HIP_CHECK(hipDeviceCanAccessPeer(&can_access_peer, src_device, dst_device));
     if (!can_access_peer) {
       std::string msg = "Skipped as peer access cannot be enabled between devices " +
-                         std::to_string(src_device) + " " + std::to_string(dst_device);
+          std::to_string(src_device) + " " + std::to_string(dst_device);
       HipTest::HIP_SKIP_TEST(msg.c_str());
       return;
     }
@@ -205,7 +206,8 @@ void Memcpy3DDeviceToDeviceShell(F memcpy_func, hipStream_t kernel_stream = null
   // Using dst_alloc width and height to set only the elements that will be copied over to
   // dst_alloc
   Iota<<<blocks, threads_per_block, 0, kernel_stream>>>(src_alloc.ptr(), src_alloc.pitch(),
-                          dst_alloc.width_logical(),dst_alloc.height(), dst_alloc.depth());
+                                                        dst_alloc.width_logical(),
+                                                        dst_alloc.height(), dst_alloc.depth());
   HIP_CHECK(hipGetLastError());
 
   HIP_CHECK(memcpy_func(dst_alloc.pitched_ptr(), make_hipPos(0, 0, 0), src_alloc.pitched_ptr(),
@@ -626,15 +628,14 @@ constexpr auto MemTypeUnified() {
 
 using DrvPtrVariant = std::variant<hipPitchedPtr, hipArray_t>;
 
-template <bool async = false>
-hipError_t DrvMemcpy3DWrapper(DrvPtrVariant dst_ptr, hipPos dst_pos, DrvPtrVariant src_ptr,
-                              hipPos src_pos, hipExtent extent, hipMemcpyKind kind,
-                              hipStream_t stream = nullptr) {
+static inline HIP_MEMCPY3D GetDrvMemcpy3DParms(DrvPtrVariant dst_ptr, hipPos dst_pos,
+                                               DrvPtrVariant src_ptr, hipPos src_pos,
+                                               hipExtent extent, hipMemcpyKind kind) {
   HIP_MEMCPY3D parms = {0};
 
   if (std::holds_alternative<hipArray_t>(dst_ptr)) {
     parms.dstMemoryType = hipMemoryTypeArray;
-    parms.dstArray = std::get<hipArray_t>(dst_ptr);  
+    parms.dstArray = std::get<hipArray_t>(dst_ptr);
   } else {
     auto ptr = std::get<hipPitchedPtr>(dst_ptr);
     parms.dstPitch = ptr.pitch;
@@ -694,6 +695,81 @@ hipError_t DrvMemcpy3DWrapper(DrvPtrVariant dst_ptr, hipPos dst_pos, DrvPtrVaria
   parms.dstY = dst_pos.y;
   parms.dstZ = dst_pos.z;
 
+  return parms;
+}
+
+static inline bool operator==(const HIP_MEMCPY3D& lhs, const HIP_MEMCPY3D& rhs) {
+  bool pos_eq = lhs.dstXInBytes == rhs.dstXInBytes && lhs.dstY == rhs.dstY &&
+      lhs.dstZ == rhs.dstZ && lhs.srcXInBytes == rhs.srcXInBytes && lhs.srcY == rhs.srcY &&
+      lhs.srcZ == rhs.srcZ;
+  bool extent_eq =
+      lhs.WidthInBytes == rhs.WidthInBytes && lhs.Height == rhs.Height && lhs.Depth == rhs.Depth;
+  bool mem_eq = true;
+  if (lhs.dstArray) {
+    mem_eq = lhs.dstArray == rhs.dstArray && lhs.dstMemoryType == rhs.dstMemoryType;
+  } else {
+    mem_eq = lhs.dstPitch == rhs.dstPitch && lhs.dstMemoryType == rhs.dstMemoryType;
+  }
+  if (lhs.srcArray) {
+    mem_eq = lhs.srcArray == rhs.srcArray && lhs.srcMemoryType == rhs.srcMemoryType;
+  } else {
+    mem_eq = lhs.srcPitch == rhs.srcPitch && lhs.srcMemoryType == rhs.srcMemoryType;
+  }
+  if (lhs.dstDevice) {
+    mem_eq = mem_eq && (lhs.dstDevice == rhs.dstDevice);
+  }
+  if (lhs.dstHost) {
+    mem_eq = mem_eq && (lhs.dstDevice == rhs.dstDevice);
+  }
+  if (lhs.srcDevice) {
+    mem_eq = mem_eq && (lhs.srcDevice == rhs.srcDevice);
+  }
+  if (lhs.srcHost) {
+    mem_eq = mem_eq && (lhs.srcHost == rhs.srcHost);
+  }
+
+  return pos_eq && extent_eq && mem_eq;
+}
+
+template <bool set_params = false>
+hipError_t DrvMemcpy3DGraphWrapper(DrvPtrVariant dst_ptr, hipPos dst_pos, DrvPtrVariant src_ptr,
+                                   hipPos src_pos, hipExtent extent, hipMemcpyKind kind,
+                                   hipCtx_t context, hipStream_t stream = nullptr) {
+  auto parms = GetDrvMemcpy3DParms(dst_ptr, dst_pos, src_ptr, src_pos, extent, kind);
+
+  hipGraph_t g = nullptr;
+  HIP_CHECK(hipGraphCreate(&g, 0));
+  hipGraphNode_t node = nullptr;
+  if constexpr (set_params) {
+    auto reversed_parms = GetDrvMemcpy3DParms(src_ptr, src_pos, dst_ptr, dst_pos, extent,
+                                              ReverseMemcpyDirection(kind));
+    HIP_CHECK(hipDrvGraphAddMemcpyNode(&node, g, nullptr, 0, &reversed_parms, context));
+    HIP_CHECK(hipDrvGraphMemcpyNodeSetParams(node, &parms));
+  } else {
+    HIP_CHECK(hipDrvGraphAddMemcpyNode(&node, g, nullptr, 0, &parms, context));
+  }
+
+  HIP_MEMCPY3D retrieved_params = {0};
+  HIP_CHECK(hipDrvGraphMemcpyNodeGetParams(node, &retrieved_params));
+  REQUIRE(parms == retrieved_params);
+
+  hipGraphExec_t graph_exec = nullptr;
+  HIP_CHECK(hipGraphInstantiate(&graph_exec, g, nullptr, nullptr, 0));
+  HIP_CHECK(hipGraphLaunch(graph_exec, hipStreamPerThread));
+  HIP_CHECK(hipStreamSynchronize(hipStreamPerThread));
+
+  HIP_CHECK(hipGraphExecDestroy(graph_exec));
+  HIP_CHECK(hipGraphDestroy(g));
+
+  return hipSuccess;
+}
+
+template <bool async = false>
+hipError_t DrvMemcpy3DWrapper(DrvPtrVariant dst_ptr, hipPos dst_pos, DrvPtrVariant src_ptr,
+                              hipPos src_pos, hipExtent extent, hipMemcpyKind kind,
+                              hipStream_t stream = nullptr) {
+  auto parms = GetDrvMemcpy3DParms(dst_ptr, dst_pos, src_ptr, src_pos, extent, kind);
+
   if constexpr (async) {
     return hipDrvMemcpy3DAsync(&parms, stream);
   } else {
diff --git a/catch/unit/graph/CMakeLists.txt b/catch/unit/graph/CMakeLists.txt
index cef5d2f5b7..82525991c7 100644
--- a/catch/unit/graph/CMakeLists.txt
+++ b/catch/unit/graph/CMakeLists.txt
@@ -151,6 +151,8 @@ set(TEST_SRC
   hipDrvGraphAddMemcpyNode.cc
   hipGraphAddMemAllocNode.cc
   hipGraphAddMemFreeNode.cc
+  hipDrvGraphMemcpyNodeGetParams.cc
+  hipDrvGraphMemcpyNodeSetParams.cc
 )
 
 add_custom_target(add_Kernel.code COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR} ${CMAKE_CURRENT_SOURCE_DIR}/add_Kernel.cpp -o ${CMAKE_CURRENT_BINARY_DIR}/../graph/add_Kernel.code -I${HIP_PATH}/include/ -I${CMAKE_CURRENT_SOURCE_DIR}/../../include --rocm-path=${ROCM_PATH})
diff --git a/catch/unit/graph/hipDrvGraphAddMemcpyNode.cc b/catch/unit/graph/hipDrvGraphAddMemcpyNode.cc
index c3bbf553db..75c20be754 100644
--- a/catch/unit/graph/hipDrvGraphAddMemcpyNode.cc
+++ b/catch/unit/graph/hipDrvGraphAddMemcpyNode.cc
@@ -17,11 +17,28 @@ OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
 THE SOFTWARE.
 */
 
+#include <functional>
+
 #include <hip_test_common.hh>
+#include <hip_test_defgroups.hh>
 #include <hip_test_checkers.hh>
+#include <memcpy3d_tests_common.hh>
+
 #include "numeric"
+#include "graph_tests_common.hh"
+
 #define XSIZE 32
 
+/**
+ * @addtogroup hipDrvGraphAddMemcpyNode hipDrvGraphAddMemcpyNode
+ * @{
+ * @ingroup GraphTest
+ * `hipDrvGraphAddMemcpyNode(hipGraphNode_t *pGraphNode, hipGraph_t graph, const
+ * hipGraphNode_t *pDependencies, size_t numDependencies, const HIP_MEMCPY3D* copyParams, hipCtx_t
+ ctx)`
+ - Creates a memcpy node and adds it to a graph
+ */
+
 /**
  * Test Description
  * ------------------------
@@ -362,3 +379,281 @@ TEST_CASE("Unit_hipDrvGraphAddMemcpyNode_MulitDevice") {
   }
 }
 #endif
+
+/**
+ * Test Description
+ * ------------------------
+ *    - Verify basic API behavior. A Memcpy node is created with parameters set according to the
+ * test run, after which the graph is run and the memcpy results are verified.
+ * The test is run for all possible memcpy directions, with both the corresponding memcpy
+ * kind and hipMemcpyDefault, as well as half page and full page allocation sizes.
+ * Test source
+ * ------------------------
+ *    - unit/graph/hipDrvGraphAddMemcpyNode.cc
+ * Test requirements
+ * ------------------------
+ *    - HIP_VERSION >= 6.0
+ */
+
+TEST_CASE("Unit_hipDrvGraphAddMemcpyNode_Positive_Basic") {
+  using namespace std::placeholders;
+
+  constexpr bool async = false;
+  HIP_CHECK(hipInit(0));
+  hipDevice_t device;
+  hipCtx_t context;
+  HIP_CHECK(hipDeviceGet(&device, 0));
+  HIP_CHECK(hipCtxCreate(&context, 0, device));
+
+  SECTION("Device to host") {
+    Memcpy3DDeviceToHostShell<async>(
+        std::bind(DrvMemcpy3DGraphWrapper<>, _1, _2, _3, _4, _5, _6, context, _7));
+  }
+
+  SECTION("Host to device") {
+    Memcpy3DHostToDeviceShell<async>(
+        std::bind(DrvMemcpy3DGraphWrapper<>, _1, _2, _3, _4, _5, _6, context, _7));
+  }
+
+  SECTION("Host to host") {
+    Memcpy3DHostToHostShell<async>(
+        std::bind(DrvMemcpy3DGraphWrapper<>, _1, _2, _3, _4, _5, _6, context, _7));
+  }
+
+  SECTION("Device to device") {
+    SECTION("Peer access enabled") {
+      Memcpy3DDeviceToDeviceShell<async, true>(
+          std::bind(DrvMemcpy3DGraphWrapper<>, _1, _2, _3, _4, _5, _6, context, _7));
+    }
+    SECTION("Peer access disabled") {
+      Memcpy3DDeviceToDeviceShell<async, false>(
+          std::bind(DrvMemcpy3DGraphWrapper<>, _1, _2, _3, _4, _5, _6, context, _7));
+    }
+  }
+
+  HIP_CHECK(hipCtxPopCurrent(&context));
+  HIP_CHECK(hipCtxDestroy(context));
+}
+
+TEST_CASE("Unit_hipDrvGraphAddMemcpyNode_Positive_Array") {
+  CHECK_IMAGE_SUPPORT
+
+  using namespace std::placeholders;
+
+  constexpr bool async = false;
+  HIP_CHECK(hipInit(0));
+  hipDevice_t device;
+  hipCtx_t context;
+  HIP_CHECK(hipDeviceGet(&device, 0));
+  HIP_CHECK(hipCtxCreate(&context, 0, device));
+
+  SECTION("Array from/to Host") {
+    DrvMemcpy3DArrayHostShell<async>(
+        std::bind(DrvMemcpy3DGraphWrapper<>, _1, _2, _3, _4, _5, _6, context, _7));
+  }
+  SECTION("Array from/to Device") {
+    DrvMemcpy3DArrayDeviceShell<async>(
+        std::bind(DrvMemcpy3DGraphWrapper<>, _1, _2, _3, _4, _5, _6, context, _7));
+  }
+
+  HIP_CHECK(hipCtxPopCurrent(&context));
+  HIP_CHECK(hipCtxDestroy(context));
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *    - Verify API behaviour with invalid arguments:
+ *        -# node is nullptr
+ *        -# graph is nullptr
+ *        -# pDependencies is nullptr when numDependencies is not zero
+ *        -# A node in pDependencies originates from a different graph
+ *        -# numDependencies is invalid
+ *        -# A node is duplicated in pDependencies
+ *        -# dst is nullptr
+ *        -# src is nullptr
+ *        -# dstPitch < width
+ *        -# srcPitch < width
+ *        -# dstPitch > max pitch
+ *        -# srcPitch > max pitch
+ *        -# WidthInBytes + dstXInBytes > dstPitch
+ *        -# WidthInBytes + srcXInBytes > srcPitch
+ *        -# dstY out of bounds
+ *        -# srcY out of bounds
+ *        -# dstZ out of bounds
+ *        -# srcZ out of bounds
+ * Test source
+ * ------------------------
+ *    - unit/graph/hipDrvGraphAddMemcpyNode.cc
+ * Test requirements
+ * ------------------------
+ *    - HIP_VERSION >= 6.0
+ */
+TEST_CASE("Unit_hipDrvGraphAddMemcpyNode_Negative_Parameters") {
+  using namespace std::placeholders;
+
+  HIP_CHECK(hipInit(0));
+  hipDevice_t device;
+  hipCtx_t context;
+  HIP_CHECK(hipDeviceGet(&device, 0));
+  HIP_CHECK(hipCtxCreate(&context, 0, device));
+
+  constexpr hipExtent extent{128 * sizeof(int), 128, 8};
+
+  constexpr auto NegativeTests = [](hipPitchedPtr dst_ptr, hipPos dst_pos, hipPitchedPtr src_ptr,
+                                    hipPos src_pos, hipExtent extent, hipMemcpyKind kind,
+                                    hipCtx_t context) {
+    hipGraph_t graph = nullptr;
+    HIP_CHECK(hipGraphCreate(&graph, 0));
+    hipGraphNode_t node = nullptr;
+
+    auto params = GetDrvMemcpy3DParms(dst_ptr, dst_pos, src_ptr, src_pos, extent, kind);
+    GraphAddNodeCommonNegativeTests(
+        std::bind(hipDrvGraphAddMemcpyNode, _1, _2, _3, _4, &params, context), graph);
+
+    SECTION("dst_ptr.ptr == nullptr") {
+      hipPitchedPtr invalid_ptr = dst_ptr;
+      invalid_ptr.ptr = nullptr;
+      auto params = GetDrvMemcpy3DParms(invalid_ptr, dst_pos, src_ptr, src_pos, extent, kind);
+      HIP_CHECK_ERROR(hipDrvGraphAddMemcpyNode(&node, graph, nullptr, 0, &params, context),
+                      hipErrorInvalidValue);
+    }
+
+    SECTION("src_ptr.ptr == nullptr") {
+      hipPitchedPtr invalid_ptr = src_ptr;
+      invalid_ptr.ptr = nullptr;
+      auto params = GetDrvMemcpy3DParms(dst_ptr, dst_pos, invalid_ptr, src_pos, extent, kind);
+      HIP_CHECK_ERROR(hipDrvGraphAddMemcpyNode(&node, graph, nullptr, 0, &params, context),
+                      hipErrorInvalidValue);
+    }
+
+    SECTION("dstPitch < width") {
+      hipPitchedPtr invalid_ptr = dst_ptr;
+      invalid_ptr.pitch = extent.width - 1;
+      auto params = GetDrvMemcpy3DParms(invalid_ptr, dst_pos, src_ptr, src_pos, extent, kind);
+      HIP_CHECK_ERROR(hipDrvGraphAddMemcpyNode(&node, graph, nullptr, 0, &params, context),
+                      hipErrorInvalidPitchValue);
+    }
+
+    SECTION("srcPitch < width") {
+      hipPitchedPtr invalid_ptr = src_ptr;
+      invalid_ptr.pitch = extent.width - 1;
+      auto params = GetDrvMemcpy3DParms(dst_ptr, dst_pos, invalid_ptr, src_pos, extent, kind);
+      HIP_CHECK_ERROR(hipDrvGraphAddMemcpyNode(&node, graph, nullptr, 0, &params, context),
+                      hipErrorInvalidPitchValue);
+    }
+
+    SECTION("dstPitch > max pitch") {
+      int attr = 0;
+      HIP_CHECK(hipDeviceGetAttribute(&attr, hipDeviceAttributeMaxPitch, 0));
+      hipPitchedPtr invalid_ptr = dst_ptr;
+      invalid_ptr.pitch = attr;
+      auto params = GetDrvMemcpy3DParms(invalid_ptr, dst_pos, src_ptr, src_pos, extent, kind);
+      HIP_CHECK_ERROR(hipDrvGraphAddMemcpyNode(&node, graph, nullptr, 0, &params, context),
+                      hipErrorInvalidValue);
+    }
+
+    SECTION("srcPitch > max pitch") {
+      int attr = 0;
+      HIP_CHECK(hipDeviceGetAttribute(&attr, hipDeviceAttributeMaxPitch, 0));
+      hipPitchedPtr invalid_ptr = src_ptr;
+      invalid_ptr.pitch = attr;
+      auto params = GetDrvMemcpy3DParms(dst_ptr, dst_pos, invalid_ptr, src_pos, extent, kind);
+      HIP_CHECK_ERROR(hipDrvGraphAddMemcpyNode(&node, graph, nullptr, 0, &params, context),
+                      hipErrorInvalidValue);
+    }
+
+    SECTION("WidthInBytes + dstXInBytes > dstPitch") {
+      hipPos invalid_pos = dst_pos;
+      invalid_pos.x = dst_ptr.pitch - extent.width + 1;
+      auto params = GetDrvMemcpy3DParms(dst_ptr, invalid_pos, src_ptr, src_pos, extent, kind);
+      HIP_CHECK_ERROR(hipDrvGraphAddMemcpyNode(&node, graph, nullptr, 0, &params, context),
+                      hipErrorInvalidValue);
+    }
+
+    SECTION("WidthInBytes + srcXInBytes > srcPitch") {
+      hipPos invalid_pos = src_pos;
+      invalid_pos.x = src_ptr.pitch - extent.width + 1;
+      auto params = GetDrvMemcpy3DParms(dst_ptr, dst_pos, src_ptr, invalid_pos, extent, kind);
+      HIP_CHECK_ERROR(hipDrvGraphAddMemcpyNode(&node, graph, nullptr, 0, &params, context),
+                      hipErrorInvalidValue);
+    }
+
+    SECTION("dstY out of bounds") {
+      hipPos invalid_pos = dst_pos;
+      invalid_pos.y = 1;
+      auto params = GetDrvMemcpy3DParms(dst_ptr, invalid_pos, src_ptr, src_pos, extent, kind);
+      HIP_CHECK_ERROR(hipDrvGraphAddMemcpyNode(&node, graph, nullptr, 0, &params, context),
+                      hipErrorInvalidValue);
+    }
+
+    SECTION("srcY out of bounds") {
+      hipPos invalid_pos = src_pos;
+      invalid_pos.y = 1;
+      auto params = GetDrvMemcpy3DParms(dst_ptr, dst_pos, src_ptr, invalid_pos, extent, kind);
+      HIP_CHECK_ERROR(hipDrvGraphAddMemcpyNode(&node, graph, nullptr, 0, &params, context),
+                      hipErrorInvalidValue);
+    }
+
+    SECTION("dstZ out of bounds") {
+      hipPos invalid_pos = dst_pos;
+      invalid_pos.z = 1;
+      auto params = GetDrvMemcpy3DParms(dst_ptr, invalid_pos, src_ptr, src_pos, extent, kind);
+      HIP_CHECK_ERROR(hipDrvGraphAddMemcpyNode(&node, graph, nullptr, 0, &params, context),
+                      hipErrorInvalidValue);
+    }
+
+    SECTION("srcZ out of bounds") {
+      hipPos invalid_pos = src_pos;
+      invalid_pos.z = 1;
+      auto params = GetDrvMemcpy3DParms(dst_ptr, dst_pos, src_ptr, invalid_pos, extent, kind);
+      HIP_CHECK_ERROR(hipDrvGraphAddMemcpyNode(&node, graph, nullptr, 0, &params, context),
+                      hipErrorInvalidValue);
+    }
+
+    HIP_CHECK(hipGraphDestroy(graph));
+  };
+
+  SECTION("Host to Device") {
+    LinearAllocGuard3D<int> device_alloc(extent);
+    LinearAllocGuard<int> host_alloc(
+        LinearAllocs::hipHostMalloc,
+        device_alloc.pitch() * device_alloc.height() * device_alloc.depth());
+    NegativeTests(device_alloc.pitched_ptr(), make_hipPos(0, 0, 0),
+                  make_hipPitchedPtr(host_alloc.ptr(), device_alloc.pitch(), device_alloc.width(),
+                                     device_alloc.height()),
+                  make_hipPos(0, 0, 0), extent, hipMemcpyHostToDevice, context);
+  }
+
+  SECTION("Device to Host") {
+    LinearAllocGuard3D<int> device_alloc(extent);
+    LinearAllocGuard<int> host_alloc(
+        LinearAllocs::hipHostMalloc,
+        device_alloc.pitch() * device_alloc.height() * device_alloc.depth());
+    NegativeTests(make_hipPitchedPtr(host_alloc.ptr(), device_alloc.pitch(), device_alloc.width(),
+                                     device_alloc.height()),
+                  make_hipPos(0, 0, 0), device_alloc.pitched_ptr(), make_hipPos(0, 0, 0), extent,
+                  hipMemcpyDeviceToHost, context);
+  }
+
+  SECTION("Host to Host") {
+    LinearAllocGuard<int> src_alloc(LinearAllocs::hipHostMalloc,
+                                    extent.width * extent.height * extent.depth);
+    LinearAllocGuard<int> dst_alloc(LinearAllocs::hipHostMalloc,
+                                    extent.width * extent.height * extent.depth);
+    NegativeTests(make_hipPitchedPtr(dst_alloc.ptr(), extent.width, extent.width, extent.height),
+                  make_hipPos(0, 0, 0),
+                  make_hipPitchedPtr(src_alloc.ptr(), extent.width, extent.width, extent.height),
+                  make_hipPos(0, 0, 0), extent, hipMemcpyHostToHost, context);
+  }
+
+  SECTION("Device to Device") {
+    LinearAllocGuard3D<int> src_alloc(extent);
+    LinearAllocGuard3D<int> dst_alloc(extent);
+    NegativeTests(dst_alloc.pitched_ptr(), make_hipPos(0, 0, 0), src_alloc.pitched_ptr(),
+                  make_hipPos(0, 0, 0), extent, hipMemcpyDeviceToDevice, context);
+  }
+
+  HIP_CHECK(hipCtxPopCurrent(&context));
+  HIP_CHECK(hipCtxDestroy(context));
+}
\ No newline at end of file
diff --git a/catch/unit/graph/hipDrvGraphMemcpyNodeGetParams.cc b/catch/unit/graph/hipDrvGraphMemcpyNodeGetParams.cc
new file mode 100644
index 0000000000..26062bbf13
--- /dev/null
+++ b/catch/unit/graph/hipDrvGraphMemcpyNodeGetParams.cc
@@ -0,0 +1,91 @@
+/*
+Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include <hip_test_defgroups.hh>
+#include <hip_test_common.hh>
+#include <memcpy3d_tests_common.hh>
+
+/**
+ * @addtogroup hipDrvGraphMemcpyNodeGetParams hipDrvGraphMemcpyNodeGetParams
+ * @{
+ * @ingroup GraphTest
+ * `hipDrvGraphMemcpyNodeGetParams(hipGraphNode_t hNode, HIP_MEMCPY3D* nodeParams)` -
+ * 	Gets a memcpy node's parameters
+ * ________________________
+ * Test cases from other APIs:
+ *  - @ref Unit_hipDrvGraphMemcpyNodeSetParams_Positive_Basic
+ */
+
+/**
+ * Test Description
+ * ------------------------
+ *    - Verify API behaviour with invalid arguments:
+ *      -# node is nullptr
+ *      -# pNodeParams is nullptr
+ *      -# node is destroyed
+ * Test source
+ * ------------------------
+ *    - unit/graph/hipDrvGraphMemcpyNodeGetParams.cc
+ * Test requirements
+ * ------------------------
+ *    - HIP_VERSION >= 6.0
+ */
+TEST_CASE("Unit_hipDrvGraphMemcpyNodeGetParams_Negative_Parameters") {
+  HIP_CHECK(hipInit(0));
+  hipDevice_t device;
+  hipCtx_t context;
+  HIP_CHECK(hipDeviceGet(&device, 0));
+  HIP_CHECK(hipCtxCreate(&context, 0, device));
+
+  constexpr hipExtent extent{128 * sizeof(int), 128, 8};
+
+  LinearAllocGuard3D<int> src_alloc(extent);
+  LinearAllocGuard3D<int> dst_alloc(extent);
+
+  auto params =
+      GetDrvMemcpy3DParms(dst_alloc.pitched_ptr(), make_hipPos(0, 0, 0), src_alloc.pitched_ptr(),
+                          make_hipPos(0, 0, 0), dst_alloc.extent(), hipMemcpyDeviceToDevice);
+
+  hipGraph_t graph = nullptr;
+  hipGraphNode_t node = nullptr;
+
+  SECTION("node == nullptr") {
+    HIP_CHECK_ERROR(hipDrvGraphMemcpyNodeGetParams(nullptr, &params), hipErrorInvalidValue);
+  }
+
+  SECTION("pNodeParams == nullptr") {
+    HIP_CHECK(hipGraphCreate(&graph, 0));
+    HIP_CHECK(hipDrvGraphAddMemcpyNode(&node, graph, nullptr, 0, &params, context));
+    HIP_CHECK_ERROR(hipDrvGraphMemcpyNodeGetParams(node, nullptr), hipErrorInvalidValue);
+    HIP_CHECK(hipGraphDestroy(graph));
+  }
+
+  SECTION("Node is destroyed") {
+    HIP_CHECK(hipGraphCreate(&graph, 0));
+    HIP_CHECK(hipDrvGraphAddMemcpyNode(&node, graph, nullptr, 0, &params, context));
+    HIP_CHECK(hipGraphDestroy(graph));
+    HIP_CHECK_ERROR(hipDrvGraphMemcpyNodeGetParams(node, &params), hipErrorInvalidValue);
+  }
+
+  HIP_CHECK(hipCtxPopCurrent(&context));
+  HIP_CHECK(hipCtxDestroy(context));
+}
diff --git a/catch/unit/graph/hipDrvGraphMemcpyNodeSetParams.cc b/catch/unit/graph/hipDrvGraphMemcpyNodeSetParams.cc
new file mode 100644
index 0000000000..b1325206e7
--- /dev/null
+++ b/catch/unit/graph/hipDrvGraphMemcpyNodeSetParams.cc
@@ -0,0 +1,314 @@
+/*
+Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include <functional>
+
+#include <hip_test_common.hh>
+#include <hip_test_defgroups.hh>
+#include <memcpy3d_tests_common.hh>
+
+/**
+ * @addtogroup hipDrvGraphMemcpyNodeSetParams hipDrvGraphMemcpyNodeSetParams
+ * @{
+ * @ingroup GraphTest
+ * `hipDrvGraphMemcpyNodeSetParams(hipGraphNode_t hNode, const HIP_MEMCPY3D* nodeParams)` - Sets a
+ * memcpy node's parameters
+ */
+
+/**
+ * Test Description
+ * ------------------------
+ *    - Verify that node parameters get updated correctly by creating a node with valid but
+ * incorrect parameters, and then setting them to the correct values after which the graph is
+ * executed and the results of the memcpy verified.
+ * The test is run for all possible memcpy directions, with both the corresponding memcpy
+ * kind and hipMemcpyDefault, as well as half page and full page allocation sizes.
+ * Test source
+ * ------------------------
+ *    - unit/graph/hipDrvGraphMemcpyNodeSetParams.cc
+ * Test requirements
+ * ------------------------
+ *    - HIP_VERSION >= 6.0
+ */
+TEST_CASE("Unit_hipDrvGraphMemcpyNodeSetParams_Positive_Basic") {
+  using namespace std::placeholders;
+
+  constexpr bool async = false;
+  HIP_CHECK(hipInit(0));
+  hipDevice_t device;
+  hipCtx_t context;
+  HIP_CHECK(hipDeviceGet(&device, 0));
+  HIP_CHECK(hipCtxCreate(&context, 0, device));
+
+  SECTION("Device to host") {
+    Memcpy3DDeviceToHostShell<async>(
+        std::bind(DrvMemcpy3DGraphWrapper<true>, _1, _2, _3, _4, _5, _6, context, _7));
+  }
+
+  SECTION("Host to device") {
+    Memcpy3DHostToDeviceShell<async>(
+        std::bind(DrvMemcpy3DGraphWrapper<true>, _1, _2, _3, _4, _5, _6, context, _7));
+  }
+
+  SECTION("Host to host") {
+    Memcpy3DHostToHostShell<async>(
+        std::bind(DrvMemcpy3DGraphWrapper<true>, _1, _2, _3, _4, _5, _6, context, _7));
+  }
+
+  SECTION("Device to device") {
+    SECTION("Peer access enabled") {
+      Memcpy3DDeviceToDeviceShell<async, true>(
+          std::bind(DrvMemcpy3DGraphWrapper<true>, _1, _2, _3, _4, _5, _6, context, _7));
+    }
+    SECTION("Peer access disabled") {
+      Memcpy3DDeviceToDeviceShell<async, false>(
+          std::bind(DrvMemcpy3DGraphWrapper<true>, _1, _2, _3, _4, _5, _6, context, _7));
+    }
+  }
+
+  HIP_CHECK(hipCtxPopCurrent(&context));
+  HIP_CHECK(hipCtxDestroy(context));
+}
+
+TEST_CASE("Unit_hipDrvGraphMemcpyNodeSetParams_Positive_Array") {
+  CHECK_IMAGE_SUPPORT
+
+  using namespace std::placeholders;
+
+  constexpr bool async = false;
+  HIP_CHECK(hipInit(0));
+  hipDevice_t device;
+  hipCtx_t context;
+  HIP_CHECK(hipDeviceGet(&device, 0));
+  HIP_CHECK(hipCtxCreate(&context, 0, device));
+
+  SECTION("Array from/to Host") {
+    DrvMemcpy3DArrayHostShell<async>(
+        std::bind(DrvMemcpy3DGraphWrapper<true>, _1, _2, _3, _4, _5, _6, context, _7));
+  }
+  SECTION("Array from/to Device") {
+    DrvMemcpy3DArrayDeviceShell<async>(
+        std::bind(DrvMemcpy3DGraphWrapper<true>, _1, _2, _3, _4, _5, _6, context, _7));
+  }
+
+  HIP_CHECK(hipCtxPopCurrent(&context));
+  HIP_CHECK(hipCtxDestroy(context));
+}
+
+
+/**
+ * Test Description
+ * ------------------------
+ *    - Verify API behaviour with invalid arguments:
+ *        -# node is nullptr
+ *        -# dst is nullptr
+ *        -# src is nullptr
+ *        -# dstPitch < width
+ *        -# srcPitch < width
+ *        -# dstPitch > max pitch
+ *        -# srcPitch > max pitch
+ *        -# WidthInBytes + dstXInBytes > dstPitch
+ *        -# WidthInBytes + srcXInBytes > srcPitch
+ *        -# dstY out of bounds
+ *        -# srcY out of bounds
+ *        -# dstZ out of bounds
+ *        -# srcZ out of bounds
+ * Test source
+ * ------------------------
+ *    - unit/graph/hipDrvGraphMemcpyNodeSetParams.cc
+ * Test requirements
+ * ------------------------
+ *    - HIP_VERSION >= 6.0
+ */
+TEST_CASE("Unit_hipDrvGraphMemcpyNodeSetParams_Negative_Parameters") {
+  using namespace std::placeholders;
+
+  HIP_CHECK(hipInit(0));
+  hipDevice_t device;
+  hipCtx_t context;
+  HIP_CHECK(hipDeviceGet(&device, 0));
+  HIP_CHECK(hipCtxCreate(&context, 0, device));
+
+  constexpr hipExtent extent{128 * sizeof(int), 128, 8};
+
+  constexpr auto NegativeTests = [](hipPitchedPtr dst_ptr, hipPos dst_pos, hipPitchedPtr src_ptr,
+                                    hipPos src_pos, hipExtent extent, hipMemcpyKind kind,
+                                    hipCtx_t context) {
+    hipGraph_t graph = nullptr;
+    HIP_CHECK(hipGraphCreate(&graph, 0));
+    hipGraphNode_t node = nullptr;
+
+    auto params = GetDrvMemcpy3DParms(dst_ptr, dst_pos, src_ptr, src_pos, extent, kind);
+    HIP_CHECK(hipDrvGraphAddMemcpyNode(&node, graph, nullptr, 0, &params, context));
+
+    SECTION("node == nullptr") {
+      HIP_CHECK_ERROR(hipDrvGraphMemcpyNodeSetParams(nullptr, &params), hipErrorInvalidValue);
+    }
+
+    SECTION("dst_ptr.ptr == nullptr") {
+      hipPitchedPtr invalid_ptr = dst_ptr;
+      invalid_ptr.ptr = nullptr;
+      auto invalid_params =
+          GetDrvMemcpy3DParms(invalid_ptr, dst_pos, src_ptr, src_pos, extent, kind);
+      HIP_CHECK_ERROR(hipDrvGraphMemcpyNodeSetParams(node, &invalid_params), hipErrorInvalidValue);
+    }
+
+    SECTION("src_ptr.ptr == nullptr") {
+      hipPitchedPtr invalid_ptr = src_ptr;
+      invalid_ptr.ptr = nullptr;
+      auto invalid_params =
+          GetDrvMemcpy3DParms(dst_ptr, dst_pos, invalid_ptr, src_pos, extent, kind);
+      HIP_CHECK_ERROR(hipDrvGraphMemcpyNodeSetParams(node, &invalid_params), hipErrorInvalidValue);
+    }
+
+    SECTION("dstPitch < width") {
+      hipPitchedPtr invalid_ptr = dst_ptr;
+      invalid_ptr.pitch = extent.width - 1;
+      auto invalid_params =
+          GetDrvMemcpy3DParms(invalid_ptr, dst_pos, src_ptr, src_pos, extent, kind);
+      HIP_CHECK_ERROR(hipDrvGraphMemcpyNodeSetParams(node, &invalid_params),
+                      hipErrorInvalidPitchValue);
+    }
+
+    SECTION("srcPitch < width") {
+      hipPitchedPtr invalid_ptr = src_ptr;
+      invalid_ptr.pitch = extent.width - 1;
+      auto invalid_params =
+          GetDrvMemcpy3DParms(dst_ptr, dst_pos, invalid_ptr, src_pos, extent, kind);
+      HIP_CHECK_ERROR(hipDrvGraphMemcpyNodeSetParams(node, &invalid_params),
+                      hipErrorInvalidPitchValue);
+    }
+
+    SECTION("dstPitch > max pitch") {
+      int attr = 0;
+      HIP_CHECK(hipDeviceGetAttribute(&attr, hipDeviceAttributeMaxPitch, 0));
+      hipPitchedPtr invalid_ptr = dst_ptr;
+      invalid_ptr.pitch = attr;
+      auto invalid_params =
+          GetDrvMemcpy3DParms(invalid_ptr, dst_pos, src_ptr, src_pos, extent, kind);
+      HIP_CHECK_ERROR(hipDrvGraphMemcpyNodeSetParams(node, &invalid_params), hipErrorInvalidValue);
+    }
+
+    SECTION("srcPitch > max pitch") {
+      int attr = 0;
+      HIP_CHECK(hipDeviceGetAttribute(&attr, hipDeviceAttributeMaxPitch, 0));
+      hipPitchedPtr invalid_ptr = src_ptr;
+      invalid_ptr.pitch = attr;
+      auto invalid_params =
+          GetDrvMemcpy3DParms(dst_ptr, dst_pos, invalid_ptr, src_pos, extent, kind);
+      HIP_CHECK_ERROR(hipDrvGraphMemcpyNodeSetParams(node, &invalid_params), hipErrorInvalidValue);
+    }
+
+    SECTION("WidthInBytes + dstXInBytes > dstPitch") {
+      hipPos invalid_pos = dst_pos;
+      invalid_pos.x = dst_ptr.pitch - extent.width + 1;
+      auto invalid_params =
+          GetDrvMemcpy3DParms(dst_ptr, invalid_pos, src_ptr, src_pos, extent, kind);
+      HIP_CHECK_ERROR(hipDrvGraphMemcpyNodeSetParams(node, &invalid_params), hipErrorInvalidValue);
+    }
+
+    SECTION("WidthInBytes + srcXInBytes > srcPitch") {
+      hipPos invalid_pos = src_pos;
+      invalid_pos.x = src_ptr.pitch - extent.width + 1;
+      auto invalid_params =
+          GetDrvMemcpy3DParms(dst_ptr, dst_pos, src_ptr, invalid_pos, extent, kind);
+      HIP_CHECK_ERROR(hipDrvGraphMemcpyNodeSetParams(node, &invalid_params), hipErrorInvalidValue);
+    }
+
+    SECTION("dstY out of bounds") {
+      hipPos invalid_pos = dst_pos;
+      invalid_pos.y = 1;
+      auto invalid_params =
+          GetDrvMemcpy3DParms(dst_ptr, invalid_pos, src_ptr, src_pos, extent, kind);
+      HIP_CHECK_ERROR(hipDrvGraphMemcpyNodeSetParams(node, &invalid_params), hipErrorInvalidValue);
+    }
+
+    SECTION("srcY out of bounds") {
+      hipPos invalid_pos = src_pos;
+      invalid_pos.y = 1;
+      auto invalid_params =
+          GetDrvMemcpy3DParms(dst_ptr, dst_pos, src_ptr, invalid_pos, extent, kind);
+      HIP_CHECK_ERROR(hipDrvGraphMemcpyNodeSetParams(node, &invalid_params), hipErrorInvalidValue);
+    }
+
+    SECTION("dstZ out of bounds") {
+      hipPos invalid_pos = dst_pos;
+      invalid_pos.z = 1;
+      auto invalid_params =
+          GetDrvMemcpy3DParms(dst_ptr, invalid_pos, src_ptr, src_pos, extent, kind);
+      HIP_CHECK_ERROR(hipDrvGraphMemcpyNodeSetParams(node, &invalid_params), hipErrorInvalidValue);
+    }
+
+    SECTION("srcZ out of bounds") {
+      hipPos invalid_pos = src_pos;
+      invalid_pos.z = 1;
+      auto invalid_params =
+          GetDrvMemcpy3DParms(dst_ptr, dst_pos, src_ptr, invalid_pos, extent, kind);
+      HIP_CHECK_ERROR(hipDrvGraphMemcpyNodeSetParams(node, &invalid_params), hipErrorInvalidValue);
+    }
+
+    HIP_CHECK(hipGraphDestroy(graph));
+  };
+
+  SECTION("Host to Device") {
+    LinearAllocGuard3D<int> device_alloc(extent);
+    LinearAllocGuard<int> host_alloc(
+        LinearAllocs::hipHostMalloc,
+        device_alloc.pitch() * device_alloc.height() * device_alloc.depth());
+    NegativeTests(device_alloc.pitched_ptr(), make_hipPos(0, 0, 0),
+                  make_hipPitchedPtr(host_alloc.ptr(), device_alloc.pitch(), device_alloc.width(),
+                                     device_alloc.height()),
+                  make_hipPos(0, 0, 0), extent, hipMemcpyHostToDevice, context);
+  }
+
+  SECTION("Device to Host") {
+    LinearAllocGuard3D<int> device_alloc(extent);
+    LinearAllocGuard<int> host_alloc(
+        LinearAllocs::hipHostMalloc,
+        device_alloc.pitch() * device_alloc.height() * device_alloc.depth());
+    NegativeTests(make_hipPitchedPtr(host_alloc.ptr(), device_alloc.pitch(), device_alloc.width(),
+                                     device_alloc.height()),
+                  make_hipPos(0, 0, 0), device_alloc.pitched_ptr(), make_hipPos(0, 0, 0), extent,
+                  hipMemcpyDeviceToHost, context);
+  }
+
+  SECTION("Host to Host") {
+    LinearAllocGuard<int> src_alloc(LinearAllocs::hipHostMalloc,
+                                    extent.width * extent.height * extent.depth);
+    LinearAllocGuard<int> dst_alloc(LinearAllocs::hipHostMalloc,
+                                    extent.width * extent.height * extent.depth);
+    NegativeTests(make_hipPitchedPtr(dst_alloc.ptr(), extent.width, extent.width, extent.height),
+                  make_hipPos(0, 0, 0),
+                  make_hipPitchedPtr(src_alloc.ptr(), extent.width, extent.width, extent.height),
+                  make_hipPos(0, 0, 0), extent, hipMemcpyHostToHost, context);
+  }
+
+  SECTION("Device to Device") {
+    LinearAllocGuard3D<int> src_alloc(extent);
+    LinearAllocGuard3D<int> dst_alloc(extent);
+    NegativeTests(dst_alloc.pitched_ptr(), make_hipPos(0, 0, 0), src_alloc.pitched_ptr(),
+                  make_hipPos(0, 0, 0), extent, hipMemcpyDeviceToDevice, context);
+  }
+
+  HIP_CHECK(hipCtxPopCurrent(&context));
+  HIP_CHECK(hipCtxDestroy(context));
+}