From 750df46bf8665a86ceddfdbe52009e6c2661a3b4 Mon Sep 17 00:00:00 2001
From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com>
Date: Thu, 24 Feb 2022 20:47:07 +0530
Subject: [PATCH] SWDEV-316467 - tests to verify
 hipGraphMemcpyNodeSetParamsToSymbol negative scenarios (#2499)

Change-Id: I38c7fe6ea060ce8f4c2ffe3ddd6ef927da2e864a

[ROCm/hip-tests commit: 9ba5219d2374f53858da4b91e1ed2bcf7b1dd1f3]
---
 .../hip-tests/catch/unit/graph/CMakeLists.txt |   1 +
 .../hipGraphMemcpyNodeSetParamsToSymbol.cc    | 107 ++++++++++++++++++
 2 files changed, 108 insertions(+)
 create mode 100644 projects/hip-tests/catch/unit/graph/hipGraphMemcpyNodeSetParamsToSymbol.cc

diff --git a/projects/hip-tests/catch/unit/graph/CMakeLists.txt b/projects/hip-tests/catch/unit/graph/CMakeLists.txt
index 34fbb28daf..aa23776294 100644
--- a/projects/hip-tests/catch/unit/graph/CMakeLists.txt
+++ b/projects/hip-tests/catch/unit/graph/CMakeLists.txt
@@ -34,6 +34,7 @@ set(TEST_SRC
   hipGraphExecHostNodeSetParams.cc
   hipGraphAddMemcpyNodeToSymbol.cc
   hipGraphExecMemsetNodeSetParams.cc
+  hipGraphMemcpyNodeSetParamsToSymbol.cc
 )
 
 hip_add_exe_to_target(NAME GraphsTest
diff --git a/projects/hip-tests/catch/unit/graph/hipGraphMemcpyNodeSetParamsToSymbol.cc b/projects/hip-tests/catch/unit/graph/hipGraphMemcpyNodeSetParamsToSymbol.cc
new file mode 100644
index 0000000000..804c502a0c
--- /dev/null
+++ b/projects/hip-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));
+}
\ No newline at end of file