AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Revert - Changes related to hipMemcpyWithStream (#1718)

Browse Source 
Reverting #1673, #1697 and #1707.
Support for hipMemcpyWithStream and memcpy optimizations, will be brought in again once issues seen with these are resolved independently.
This commit is contained in:
Rahul Garg
2019-12-05 20:21:53 -08:00
committed by Maneesh Gupta
parent 496f0c53b3
commit e53fc316f1
4 changed files with 133 additions and 455 deletions
Download Patch File Download Diff File Expand all files Collapse all files
include/hip/hcc_detail/hip_runtime_api.h
+2 -6
View File
@@ -1277,7 +1277,7 @@ hipError_t hipMallocPitch(void** ptr, size_t* pitch, size_t width, size_t height
* @param[in] height Requested pitched allocation height
*
* If size is 0, no memory is allocated, *ptr returns nullptr, and hipSuccess is returned.
* The intended usage of pitch is as a separate parameter of the allocation, used to compute addresses within the 2D array.
* The intended usage of pitch is as a separate parameter of the allocation, used to compute addresses within the 2D array.
* Given the row and column of an array element of type T, the address is computed as:
* T* pElement = (T*)((char*)BaseAddress + Row * Pitch) + Column;
*
@@ -1361,10 +1361,6 @@ hipError_t hipHostFree(void* ptr);
*/
hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind);
// TODO: Add description
hipError_t hipMemcpyWithStream(void* dst, const void* src, size_t sizeBytes,
hipMemcpyKind kind, hipStream_t stream);
/**
* @brief Copy data from Host to Device
*
@@ -1502,7 +1498,7 @@ hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* symbolName,
#else
hipError_t hipModuleGetGlobal(void**, size_t*, hipModule_t, const char*);
#ifdef __cplusplus //Start : Not supported in gcc
#ifdef __cplusplus //Start : Not supported in gcc
namespace hip_impl {
inline
__attribute__((visibility("hidden")))
include/hip/nvcc_detail/hip_runtime_api.h
-10
View File
@@ -610,16 +610,6 @@ inline static hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes,
cudaMemcpy(dst, src, sizeBytes, hipMemcpyKindToCudaMemcpyKind(copyKind)));
}
inline hipError_t hipMemcpyWithStream(void* dst, const void* src,
size_t sizeBytes, hipMemcpyKind copyKind,
hipStream_t stream) {
cudaError_t error = cudaMemcpyAsync(dst, src, sizeBytes,
hipMemcpyKindToCudaMemcpyKind(copyKind),
stream);
if (error != cudaSuccess) return hipCUDAErrorTohipError(error);
return hipCUDAErrorTohipError(cudaStreamSynchronize(stream));
}
inline static hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes,
hipMemcpyKind copyKind, hipStream_t stream __dparm(0)) {
src/hip_memory.cpp
+131 -383
View File
@@ -27,7 +27,6 @@ THE SOFTWARE.
#include "hip_hcc_internal.h"
#include "trace_helper.h"
#include <functional>
#include <fstream>
__device__ char __hip_device_heap[__HIP_SIZE_OF_HEAP];
@@ -36,369 +35,23 @@ __device__ uint32_t __hip_device_page_flag[__HIP_NUM_PAGES];
// Internal HIP APIS:
namespace hip_internal {
namespace {
inline
const char* hsa_to_string(hsa_status_t err) noexcept
{
const char* r{};
hipError_t memcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind,
hipStream_t stream) {
hipError_t e = hipSuccess;
if (hsa_status_string(err, &r) != HSA_STATUS_SUCCESS) return r;
return "Unknown.";
}
template<std::size_t m, std::size_t n>
inline
void throwing_result_check(hsa_status_t res, const char (&file)[m],
const char (&function)[n], int line) {
if (res == HSA_STATUS_SUCCESS) return;
if (res == HSA_STATUS_INFO_BREAK) return;
throw std::runtime_error{"Failed in file " + (file +
(", in function \"" + (function +
("\", on line " + std::to_string(line))))) +
", with error: " + hsa_to_string(res)};
}
inline
hsa_agent_t cpu_agent() noexcept {
static hsa_agent_t cpu{[]() {
hsa_agent_t r{};
throwing_result_check(
hsa_iterate_agents([](hsa_agent_t x, void* pr) {
hsa_device_type_t t{};
hsa_agent_get_info(x, HSA_AGENT_INFO_DEVICE, &t);
if (t != HSA_DEVICE_TYPE_CPU) return HSA_STATUS_SUCCESS;
*static_cast<hsa_agent_t *>(pr) = x;
return HSA_STATUS_INFO_BREAK;
}, &r), __FILE__, __func__, __LINE__);
return r;
}()};
return cpu;
}
inline
hsa_device_type_t type(hsa_agent_t x) noexcept
{
hsa_device_type_t r{};
throwing_result_check(hsa_agent_get_info(x, HSA_AGENT_INFO_DEVICE, &r),
__FILE__, __func__, __LINE__);
return r;
}
const auto is_large_BAR{[](){
std::unique_ptr<void, void (*)(void*)> hsa{
(throwing_result_check(hsa_init(), __FILE__, __func__, __LINE__),
nullptr),
[](void*) { hsa_shut_down(); }};
bool r{true};
throwing_result_check(hsa_iterate_agents([](hsa_agent_t x, void* pr) {
if (x.handle == cpu_agent().handle) return HSA_STATUS_SUCCESS;
throwing_result_check(
hsa_agent_iterate_regions(x, [](hsa_region_t y, void* p) {
hsa_region_segment_t seg{};
throwing_result_check(
hsa_region_get_info(y, HSA_REGION_INFO_SEGMENT, &seg),
__FILE__, __func__, __LINE__);
if (seg != HSA_REGION_SEGMENT_GLOBAL) {
return HSA_STATUS_SUCCESS;
}
uint32_t flags{};
throwing_result_check(hsa_region_get_info(
y, HSA_REGION_INFO_GLOBAL_FLAGS, &flags),
__FILE__, __func__, __LINE__);
if (flags & HSA_REGION_GLOBAL_FLAG_COARSE_GRAINED) {
hsa_amd_memory_pool_access_t tmp{};
throwing_result_check(
hsa_amd_agent_memory_pool_get_info(
cpu_agent(),
hsa_amd_memory_pool_t{y.handle},
HSA_AMD_AGENT_MEMORY_POOL_INFO_ACCESS,
&tmp),
__FILE__, __func__, __LINE__);
*static_cast<bool*>(p) &=
tmp != HSA_AMD_MEMORY_POOL_ACCESS_NEVER_ALLOWED;
}
return HSA_STATUS_SUCCESS;
}, pr), __FILE__, __func__, __LINE__);
return HSA_STATUS_SUCCESS;
}, &r), __FILE__, __func__, __LINE__);
return r;
}()};
inline
hsa_amd_pointer_info_t info(const void* p) noexcept
{
hsa_amd_pointer_info_t r{sizeof(hsa_amd_pointer_info_t)};
throwing_result_check(
hsa_amd_pointer_info(
const_cast<void*>(p), &r, nullptr, nullptr, nullptr),
__FILE__, __func__, __LINE__);
r.size = is_large_BAR || (type(r.agentOwner) == HSA_DEVICE_TYPE_CPU) ?
UINT32_MAX : sizeof(hsa_amd_pointer_info_t);
return r;
}
constexpr size_t staging_sz{4 * 1024 * 1024}; // 2 Pages.
thread_local const std::unique_ptr<void, void (*)(void *)> staging_buffer{
[]() {
hsa_region_t r{};
throwing_result_check(hsa_agent_iterate_regions(
cpu_agent(), [](hsa_region_t x, void *p) {
hsa_region_segment_t seg{};
throwing_result_check(
hsa_region_get_info(x, HSA_REGION_INFO_SEGMENT, &seg),
__FILE__, __func__, __LINE__);
if (seg != HSA_REGION_SEGMENT_GLOBAL) return HSA_STATUS_SUCCESS;
uint32_t flags{};
throwing_result_check(hsa_region_get_info(
x, HSA_REGION_INFO_GLOBAL_FLAGS, &flags),
__FILE__, __func__, __LINE__);
if (flags & HSA_REGION_GLOBAL_FLAG_COARSE_GRAINED) {
*static_cast<hsa_region_t *>(p) = x;
return HSA_STATUS_INFO_BREAK;
}
return HSA_STATUS_SUCCESS;
}, &r), __FILE__, __func__, __LINE__);
void *tp{};
throwing_result_check(hsa_memory_allocate(r, staging_sz, &tp),
__FILE__, __func__, __LINE__);
return tp;
}(),
[](void *ptr) { hsa_memory_free(ptr); }};
thread_local hsa_signal_t copy_signal{[]() {
hsa_agent_t cpu{cpu_agent()};
hsa_signal_t sgn{};
throwing_result_check(hsa_signal_create(1, 1, &cpu, &sgn),
__FILE__, __func__, __LINE__);
return sgn;
}()};
} // Unnamed namespace.
inline
void do_copy(void* __restrict dst, const void* __restrict src, std::size_t n,
hsa_agent_t da, hsa_agent_t sa) {
hsa_signal_silent_store_relaxed(copy_signal, 1);
throwing_result_check(
hsa_amd_memory_async_copy(dst, da, src, sa, n, 0, nullptr, copy_signal),
__FILE__, __func__, __LINE__);
while (hsa_signal_wait_relaxed(copy_signal, HSA_SIGNAL_CONDITION_EQ, 0,
UINT64_MAX, HSA_WAIT_STATE_ACTIVE));
}
inline
void do_std_memcpy(
void* __restrict dst, const void* __restrict src, std::size_t n) {
std::memcpy(dst, src, n);
return std::atomic_thread_fence(std::memory_order_seq_cst);
}
inline
void d2h_copy(void* __restrict dst, const void* __restrict src, size_t n,
hsa_amd_pointer_info_t si) {
if (si.size == UINT32_MAX) return do_std_memcpy(dst, src, n);
const auto di{info(dst)};
if (di.type == HSA_EXT_POINTER_TYPE_LOCKED) {
dst = static_cast<char*>(di.agentBaseAddress) +
(static_cast<char*>(dst) -
static_cast<char*>(di.hostBaseAddress));
do_copy(dst, src, n, si.agentOwner, si.agentOwner);
}
else if (n <= staging_sz) {
do_copy(staging_buffer.get(), src, n, si.agentOwner, si.agentOwner);
std::memcpy(dst, staging_buffer.get(), n);
}
else {
std::unique_ptr<void, void (*)(void*)> lck{
dst, [](void* p) { hsa_amd_memory_unlock(p); }};
throwing_result_check(hsa_amd_memory_lock(dst, n, &si.agentOwner, 1,
const_cast<void**>(&dst)),
__FILE__, __func__, __LINE__);
do_copy(dst, src, n, si.agentOwner, si.agentOwner);
}
}
inline
void h2d_copy(void* __restrict dst, const void* __restrict src, size_t n,
hsa_amd_pointer_info_t di) {
if (di.size == UINT32_MAX) return do_std_memcpy(dst, src, n);
const auto si{info(const_cast<void*>(src))};
if (si.type == HSA_EXT_POINTER_TYPE_LOCKED) {
src = static_cast<char*>(si.agentBaseAddress) +
(static_cast<const char*>(src) -
static_cast<char*>(si.hostBaseAddress));
do_copy(dst, src, n, di.agentOwner, di.agentOwner);
}
else if (n <= staging_sz) {
std::memcpy(staging_buffer.get(), src, n);
do_copy(dst, staging_buffer.get(), n, di.agentOwner, di.agentOwner);
}
else {
std::unique_ptr<void, void (*)(void*)> lck{
const_cast<void*>(src), [](void* p) { hsa_amd_memory_unlock(p); }};
throwing_result_check(hsa_amd_memory_lock(const_cast<void*>(src), n,
&di.agentOwner, 1,
const_cast<void**>(&src)),
__FILE__, __func__, __LINE__);
do_copy(dst, src, n, di.agentOwner, di.agentOwner);
}
}
inline
void generic_copy(void* __restrict dst, const void* __restrict src, size_t n,
hsa_amd_pointer_info_t di, hsa_amd_pointer_info_t si) {
if (di.size == UINT32_MAX && si.size == UINT32_MAX) {
return do_std_memcpy(dst, src, n);
}
std::unique_ptr<void, void (*)(void*)> lck0{
nullptr, [](void* p) { hsa_amd_memory_unlock(p); }};
std::unique_ptr<void, void (*)(void*)> lck1{nullptr, lck0.get_deleter()};
switch (si.type) {
case HSA_EXT_POINTER_TYPE_HSA:
if (di.type == HSA_EXT_POINTER_TYPE_HSA) {
hsa_memory_copy(dst, src, n);
return; // TODO: do_copy(dst, src, n, di.agentOwner, si.agentOwner);
}
if (di.type == HSA_EXT_POINTER_TYPE_UNKNOWN ||
di.type == HSA_EXT_POINTER_TYPE_LOCKED) {
return d2h_copy(dst, src, n, si);
}
break;
case HSA_EXT_POINTER_TYPE_LOCKED:
if (di.type == HSA_EXT_POINTER_TYPE_UNKNOWN) {
std::memcpy(dst, si.hostBaseAddress, n);
return;
}
if (di.type == HSA_EXT_POINTER_TYPE_LOCKED) {
std::memcpy(di.hostBaseAddress, si.hostBaseAddress, n);
return;
}
src = si.agentBaseAddress;
si.agentOwner = di.agentOwner;
break;
case HSA_EXT_POINTER_TYPE_UNKNOWN:
if (di.type == HSA_EXT_POINTER_TYPE_UNKNOWN) {
std::memcpy(dst, src, n);
return;
}
if (di.type == HSA_EXT_POINTER_TYPE_LOCKED) {
std::memcpy(di.hostBaseAddress, src, n);
return;
}
return h2d_copy(dst, src, n, di);
default: do_copy(dst, src, n, di.agentOwner, si.agentOwner); break;
}
}
inline
void memcpy_impl(void* __restrict dst, const void* __restrict src, size_t n,
hipMemcpyKind k) noexcept {
switch (k) {
case hipMemcpyHostToHost: std::memcpy(dst, src, n); break;
case hipMemcpyHostToDevice:
return is_large_BAR ? do_std_memcpy(dst, src, n)
: h2d_copy(dst, src, n, info(dst));
case hipMemcpyDeviceToHost:
return is_large_BAR ? do_std_memcpy(dst, src, n)
: d2h_copy(dst, src, n, info(src));
case hipMemcpyDeviceToDevice:
return do_copy(dst, src, n, info(dst).agentOwner,
info(const_cast<void*>(src)).agentOwner);
default: return generic_copy(dst, src, n, info(dst), info(src));
}
}
hipError_t memcpyAsync(void* dst, const void* src, size_t sizeBytes,
hipMemcpyKind kind, hipStream_t stream) {
if (sizeBytes == 0) return hipSuccess;
// Return success if number of bytes to copy is 0
if (sizeBytes == 0) return e;
if (!dst || !src) return hipErrorInvalidValue;
if (!(stream = ihipSyncAndResolveStream(stream))) {
return hipErrorInvalidValue;
}
try {
stream = ihipSyncAndResolveStream(stream);
if (!stream) return hipErrorInvalidValue;
stream->locked_copyAsync(dst, src, sizeBytes, kind);
}
catch (const ihipException& ex) {
return ex._code;
}
catch (const std::exception& ex) {
std::cerr << ex.what() << std::endl;
throw;
}
catch (...) {
return hipErrorUnknown;
}
return hipSuccess;
}
hipError_t memcpySync(void* dst, const void* src, size_t sizeBytes,
hipMemcpyKind kind, hipStream_t stream) {
if (sizeBytes == 0) return hipSuccess;
if (!dst || !src) return hipErrorInvalidValue;
try {
stream = ihipSyncAndResolveStream(stream);
if (!stream) return hipErrorInvalidValue;
LockedAccessor_StreamCrit_t cs{stream->criticalData()};
cs->_av.wait();
memcpy_impl(dst, src, sizeBytes, kind);
}
catch (const ihipException& ex) {
return ex._code;
}
catch (const std::exception& ex) {
std::cerr << ex.what() << std::endl;
throw;
catch (ihipException& ex) {
e = ex._code;
}
catch (...) {
return hipErrorUnknown;
@@ -1243,8 +896,20 @@ hipError_t hipMemcpyToSymbol(void* dst, const void* src, size_t count,
tprintf(DB_MEM, " symbol '%s' resolved to address:%p\n", symbol_name, dst);
return ihipLogStatus(
hipMemcpy(static_cast<char*>(dst) + offset, src, count, kind));
if (dst == nullptr) {
return ihipLogStatus(hipErrorInvalidSymbol);
}
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
if (kind == hipMemcpyHostToDevice || kind == hipMemcpyDefault ||
kind == hipMemcpyDeviceToDevice || kind == hipMemcpyHostToHost) {
stream->locked_copySync((char*)dst+offset, (void*)src, count, kind, false);
} else {
return ihipLogStatus(hipErrorInvalidValue);
}
return ihipLogStatus(hipSuccess);
}
hipError_t hipMemcpyFromSymbol(void* dst, const void* src, size_t count,
@@ -1255,8 +920,20 @@ hipError_t hipMemcpyFromSymbol(void* dst, const void* src, size_t count,
tprintf(DB_MEM, " symbol '%s' resolved to address:%p\n", symbol_name, dst);
return ihipLogStatus(
hipMemcpy(dst, static_cast<const char*>(src) + offset, count, kind));
if (dst == nullptr) {
return ihipLogStatus(hipErrorInvalidSymbol);
}
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
if (kind == hipMemcpyDefault || kind == hipMemcpyDeviceToHost ||
kind == hipMemcpyDeviceToDevice || kind == hipMemcpyHostToHost) {
stream->locked_copySync((void*)dst, (char*)src+offset, count, kind, false);
} else {
return ihipLogStatus(hipErrorInvalidValue);
}
return ihipLogStatus(hipSuccess);
}
@@ -1318,49 +995,120 @@ hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* src, size_t count,
hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind) {
HIP_INIT_SPECIAL_API(hipMemcpy, (TRACE_MCMD), dst, src, sizeBytes, kind);
return ihipLogStatus(hip_internal::memcpySync(dst, src, sizeBytes, kind,
hipStreamNull));
hipError_t e = hipSuccess;
// Return success if number of bytes to copy is 0
if (sizeBytes == 0) return ihipLogStatus(e);
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
hc::completion_future marker;
if(dst==NULL || src==NULL)
{
e=hipErrorInvalidValue;
return ihipLogStatus(e);
}
try {
stream->locked_copySync(dst, src, sizeBytes, kind);
} catch (ihipException& ex) {
e = ex._code;
}
return ihipLogStatus(e);
}
hipError_t hipMemcpyHtoD(hipDeviceptr_t dst, void* src, size_t sizeBytes) {
HIP_INIT_SPECIAL_API(hipMemcpyHtoD, (TRACE_MCMD), dst, src, sizeBytes);
return ihipLogStatus(hip_internal::memcpySync(dst, src, sizeBytes,
hipMemcpyHostToDevice,
hipStreamNull));
hipError_t e = hipSuccess;
if (sizeBytes == 0) return ihipLogStatus(e);
if(dst==NULL || src==NULL){
return ihipLogStatus(hipErrorInvalidValue);
}
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
hc::completion_future marker;
try {
stream->locked_copySync((void*)dst, (void*)src, sizeBytes, hipMemcpyHostToDevice, false);
} catch (ihipException& ex) {
e = ex._code;
}
return ihipLogStatus(e);
}
hipError_t hipMemcpyDtoH(void* dst, hipDeviceptr_t src, size_t sizeBytes) {
HIP_INIT_SPECIAL_API(hipMemcpyDtoH, (TRACE_MCMD), dst, src, sizeBytes);
return ihipLogStatus(hip_internal::memcpySync(dst, src, sizeBytes,
hipMemcpyDeviceToHost,
hipStreamNull));
hipError_t e = hipSuccess;
if (sizeBytes == 0) return ihipLogStatus(e);
if(dst==NULL || src==NULL){
return ihipLogStatus(hipErrorInvalidValue);
}
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
hc::completion_future marker;
try {
stream->locked_copySync((void*)dst, (void*)src, sizeBytes, hipMemcpyDeviceToHost, false);
} catch (ihipException& ex) {
e = ex._code;
}
return ihipLogStatus(e);
}
hipError_t hipMemcpyDtoD(hipDeviceptr_t dst, hipDeviceptr_t src, size_t sizeBytes) {
HIP_INIT_SPECIAL_API(hipMemcpyDtoD, (TRACE_MCMD), dst, src, sizeBytes);
return ihipLogStatus(hip_internal::memcpySync(dst, src, sizeBytes,
hipMemcpyDeviceToDevice,
hipStreamNull));
hipError_t e = hipSuccess;
if (sizeBytes == 0) return ihipLogStatus(e);
if(dst==NULL || src==NULL){
return ihipLogStatus(hipErrorInvalidValue);
}
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
hc::completion_future marker;
try {
stream->locked_copySync((void*)dst, (void*)src, sizeBytes, hipMemcpyDeviceToDevice, false);
} catch (ihipException& ex) {
e = ex._code;
}
return ihipLogStatus(e);
}
hipError_t hipMemcpyHtoH(void* dst, void* src, size_t sizeBytes) {
HIP_INIT_SPECIAL_API(hipMemcpyHtoH, (TRACE_MCMD), dst, src, sizeBytes);
return ihipLogStatus(hip_internal::memcpySync(dst, src, sizeBytes,
hipMemcpyHostToHost,
hipStreamNull));
}
hipError_t e = hipSuccess;
if (sizeBytes == 0) return ihipLogStatus(e);
hipError_t hipMemcpyWithStream(void* dst, const void* src, size_t sizeBytes,
hipMemcpyKind kind, hipStream_t stream) {
HIP_INIT_SPECIAL_API(hipMemcpyWithStream, (TRACE_MCMD), dst, src, sizeBytes,
kind, stream);
if(dst==NULL || src==NULL){
return ihipLogStatus(hipErrorInvalidValue);
}
return ihipLogStatus(hip_internal::memcpySync(dst, src, sizeBytes, kind,
stream));
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
hc::completion_future marker;
try {
stream->locked_copySync((void*)dst, (void*)src, sizeBytes, hipMemcpyHostToHost, false);
} catch (ihipException& ex) {
e = ex._code;
}
return ihipLogStatus(e);
}
hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind,
tests/src/runtimeApi/memory/hipMemcpyWithStream.cpp
-56
View File
@@ -1,56 +0,0 @@
/*
Copyright (c) 2015-present 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.
*/
/*
* Conformance test for checking functionality of
* hipError_t hipMemcpyPeer(void* dst, int dstDeviceId, const void* src, int srcDeviceId, size_t
* sizeBytes);
*/
/* HIT_START
* BUILD: %t %s ../../test_common.cpp
* TEST: %t
* HIT_END
*/
#include "test_common.h"
int main() {
size_t Nbytes = N * sizeof(int);
int numDevices = 0;
int *A_d, *B_d, *C_d;
int *A_h, *B_h, *C_h;
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
hipStream_t stream;
HIPCHECK(hipStreamCreate(&stream));
HIPCHECK(hipMemcpyWithStream(A_d, A_h, Nbytes, hipMemcpyHostToDevice, stream));
HIPCHECK(hipMemcpyWithStream(B_d, B_h, Nbytes, hipMemcpyHostToDevice,stream));
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, stream,
static_cast<const int*>(A_d), static_cast<const int*>(B_d), C_d, N);
HIPCHECK(hipStreamSynchronize(stream));
HIPCHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
HipTest::checkVectorADD(A_h, B_h, C_h, N);
HIPCHECK(hipStreamDestroy(stream));
passed();
}