SWDEV-327960, SWDEV-316901 - Adding Graph checks

hipGraphExecMemcpyNodeSetParamsToSymbol, fix for functional cases
failing as well

Change-Id: I9f263778f6b24e3ec35a78231a17e01c0d44f23a


[ROCm/clr commit: 63f23ec2b0]

projects/clr/hipamd/src/hip_graph_internal.hpp

@@ -788,32 +788,26 @@ class hipGraphMemcpyNodeToSymbol : public hipGraphMemcpyNode1D {
   hipError_t SetParams(const void* symbol, const void* src, size_t count, size_t offset,
                        hipMemcpyKind kind) {
 
-    size_t zeroOffset = 0;
-    amd::Memory* srcMemory = getMemoryObject(src, zeroOffset);
-    amd::Memory* dstMemory = getMemoryObject(symbol, zeroOffset);
-    hipMemoryType srcMemoryType =
-        amd::MemObjMap::FindMemObj(srcMemory) ? hipMemoryTypeDevice : hipMemoryTypeHost;
-    hipMemoryType dstMemoryType =
-        amd::MemObjMap::FindMemObj(dstMemory) ? hipMemoryTypeDevice : hipMemoryTypeHost;
-
-    // Return error if sizeBytes passed to memcpy is more than the actual size allocated
-    if ((dstMemory && count > (dstMemory->getSize() - offset)) ||
-        (srcMemory && count > (srcMemory->getSize() - offset))) {
-      return hipErrorInvalidValue;
-    }
-    // check the kind with memory types
-    if (std::get<0>(hip::getMemoryType(kind)) != srcMemoryType)
-      return hipErrorInvalidValue;
-    if (std::get<1>(hip::getMemoryType(kind)) != dstMemoryType)
-        return hipErrorInvalidValue;
-
     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 src is also a symbol (cuda negative test case)
+    hipError_t status = ihipMemcpySymbol_validate(src, 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* srcMemory = getMemoryObject(src, dOffset);
+    if(srcMemory == nullptr && kind != hipMemcpyHostToDevice) {
+      return hipErrorInvalidValue;
+    } else if (srcMemory != nullptr && kind != hipMemcpyDeviceToDevice) {
+      return hipErrorInvalidValue;
+    } else if (kind == hipMemcpyHostToHost || kind == hipMemcpyDeviceToHost) {
+      return hipErrorInvalidValue;
+    }
     symbol_ = symbol;
     src_ = src;
     count_ = count;