From 798a304dbdea92cc635006422b90d4201e3ccf6c Mon Sep 17 00:00:00 2001
From: pghafari <pghafari@amd.com>
Date: Mon, 4 Apr 2022 23:29:56 -0400
Subject: [PATCH] SWDEV-316292 - adding Graph checks

hipGraphExecMemcpyNodeSetParamsFromSymbol negative test cases

Change-Id: I7b6c496a8bca831dc90d1d427273a2a4fea96587
---
 hipamd/src/hip_graph_internal.hpp | 19 +++++++++++++++++--
 1 file changed, 17 insertions(+), 2 deletions(-)

diff --git a/hipamd/src/hip_graph_internal.hpp b/hipamd/src/hip_graph_internal.hpp
index eeb5c05cf5..325b3dd5b1 100644
--- a/hipamd/src/hip_graph_internal.hpp
+++ b/hipamd/src/hip_graph_internal.hpp
@@ -699,11 +699,26 @@ class hipGraphMemcpyNodeFromSymbol : public hipGraphMemcpyNode1D {
                        hipMemcpyKind kind) {
     size_t sym_size = 0;
     hipDeviceptr_t device_ptr = nullptr;
-
-    hipError_t status = ihipMemcpySymbol_validate(symbol, count, offset, sym_size, device_ptr);
+    //check to see if dst is also a symbol (cuda negative test case)
+    hipError_t status = ihipMemcpySymbol_validate(dst, count, offset, sym_size, device_ptr);
+    if (status == hipSuccess) {
+      return hipErrorInvalidValue;
+    }
+    status = ihipMemcpySymbol_validate(symbol, count, offset, sym_size, device_ptr);
     if (status != hipSuccess) {
       return status;
     }
+
+    size_t dOffset = 0;
+    amd::Memory* dstMemory = getMemoryObject(dst, dOffset);
+    if( dstMemory == nullptr && kind != hipMemcpyHostToDevice) {
+      return hipErrorInvalidMemcpyDirection;
+    } else if ( dstMemory != nullptr && kind != hipMemcpyDeviceToDevice) {
+      return hipErrorInvalidMemcpyDirection;
+    } else if ( kind == hipMemcpyHostToHost || kind == hipMemcpyDeviceToHost) {
+      return hipErrorInvalidMemcpyDirection;
+    }
+
     dst_ = dst;
     symbol_ = symbol;
     count_ = count;