diff --git a/src/hip_memory.cpp b/src/hip_memory.cpp index 4a126532f4..3c9b1616af 100644 --- a/src/hip_memory.cpp +++ b/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); } } diff --git a/tests/src/runtimeApi/memory/hipMemcpyNegetiveTests.cpp b/tests/src/runtimeApi/memory/hipMemcpyNegetiveTests.cpp new file mode 100644 index 0000000000..692d14cec7 --- /dev/null +++ b/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(); +}