Fix for segfault seen if invalid kind is passed to hipMemcpy (#1937)

Fixes SWDEV-224941

src/hip_memory.cpp

@@ -320,20 +320,34 @@ void generic_copy(void* __restrict dst, const void* __restrict src, size_t n,
 inline
 void memcpy_impl(void* __restrict dst, const void* __restrict src, size_t n,
                  hipMemcpyKind k) {
+    auto si{info(src)};
+    auto di{info(dst)};
+
+    if (!is_large_BAR){
+       // Pointer info takes presidence over hipMemcpyKind
+       // if there is mismatch b/w Memcpy kind and dst/src pointer
+       // E.g. dst(host pointer),src(device pointer) and hipMemcpyKind set as hipMemcpyHostToDevice
+       if (di.size == is_cpu_owned && si.size == is_cpu_owned)
+          k = hipMemcpyHostToHost;
+       else if (si.size == is_cpu_owned && di.size != is_cpu_owned)
+          k = hipMemcpyHostToDevice;
+       else if (di.size == is_cpu_owned && si.size != is_cpu_owned)
+          k = hipMemcpyDeviceToHost;
+       else
+          k = hipMemcpyDeviceToDevice;
+    }
     switch (k) {
     case hipMemcpyHostToHost: std::memcpy(dst, src, n); break;
-    case hipMemcpyHostToDevice: return h2d_copy(dst, src, n, info(dst));
-    case hipMemcpyDeviceToHost: return d2h_copy(dst, src, n, info(src));
+    case hipMemcpyHostToDevice: return h2d_copy(dst, src, n, di);
+    case hipMemcpyDeviceToHost: return d2h_copy(dst, src, n, si);
     case hipMemcpyDeviceToDevice: {
-        const auto di{info(dst)};
-        const auto si{info(src)};
         throwing_result_check(hsa_amd_agents_allow_access(1u, &si.agentOwner,
                                                           nullptr,
                                                           di.agentBaseAddress),
                               __FILE__, __func__, __LINE__);
         return do_copy(dst, src, n, di.agentOwner, si.agentOwner);
     }
-    default: return generic_copy(dst, src, n, info(dst), info(src));
+    default: return generic_copy(dst, src, n, di, si);
     }
 }
 

tests/src/runtimeApi/memory/hipMemcpyNegetiveTests.cpp

@@ -0,0 +1,53 @@
+/*
+ * Copyright (c) 2019-2020 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, INNCLUDING 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 ANNY 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.
+ * */
+
+/* HIT_START
+ * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
+ * TEST: %t
+ * HIT_END
+ */
+
+
+#include "test_common.h"
+
+int main() {
+    int* A;
+    int* Ad;
+    int* Bd;
+    
+    // Allocation
+    HIPCHECK(hipMalloc((void**)&Ad, sizeof(int)));
+    HIPCHECK(hipMalloc((void**)&Bd, sizeof(int)));
+    HIPCHECK(hipHostMalloc((void**)&A,sizeof(int)));
+
+    // Kind should be ignored and test should pass even for incorrect kind
+    HIPCHECK(hipMemcpy(Ad, A, sizeof(int), hipMemcpyDeviceToHost));
+    HIPCHECK(hipMemcpy(A,  Ad, sizeof(int), hipMemcpyHostToDevice));
+    HIPCHECK(hipMemcpy(Ad, Bd, sizeof(int), hipMemcpyHostToHost));
+    HIPCHECK(hipMemcpy(A,  A, sizeof(int), hipMemcpyDeviceToDevice));
+    
+    // nullptr passed as source or destination pointer
+    HIPASSERT(hipSuccess != hipMemcpy(nullptr, A, sizeof(int), hipMemcpyHostToDevice));
+    HIPASSERT(hipSuccess != hipMemcpy(Ad, nullptr, sizeof(int), hipMemcpyHostToDevice));
+    
+    HIPCHECK(hipFree(Ad));
+    HIPCHECK(hipFree(Bd));
+    HIPCHECK(hipFree(A));
+    passed();   
+}