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

SWDEV-356299 - Fix hipPeerToPeer_simple Tests (#2953)

Browse Source 
- Add hipDeviceSynchronize after hipMemset with device memory, as hipMemset is asynchronous in this case

Change-Id: Ia55d106931984fed450be29c1b745c8069151f76
This commit is contained in:
ROCm CI Service Account
2022-09-21 14:17:35 +05:30
committed by GitHub
parent 046a7a193e
commit 6e021f1adc
1 changed files with 10 additions and 1 deletions
Download Patch File Download Diff File Expand all files Collapse all files
tests/src/p2p/hipPeerToPeer_simple.cpp
+10 -1
View File
@@ -159,12 +159,15 @@ void enablePeerFirst(bool useAsyncCopy) {
HIPCHECK(hipSetDevice(g_currentDevice));
HIPCHECK(hipMalloc(&A_d0, Nbytes));
HIPCHECK(hipMemset(A_d0, memsetval, Nbytes));
// hipDeviceSynchronize as hipMemset is asynchronous when destination memory is device memory
HIPCHECK(hipDeviceSynchronize());
// allocate and initialize memory on peer device
HIPCHECK(hipSetDevice(g_peerDevice));
HIPCHECK(hipMalloc(&A_d1, Nbytes));
HIPCHECK(hipMemset(A_d1, 0x13, Nbytes));
// hipDeviceSynchronize as hipMemset is asynchronous when destination memory is device memory
HIPCHECK(hipDeviceSynchronize());
// Device0 push to device1, using P2P:
// NOTE : if p_mirrorPeers=0 and p_memcpyWithPeer=1, then peer device does not have mapping for
@@ -216,11 +219,15 @@ void allocMemoryFirst(bool useAsyncCopy) {
HIPCHECK(hipSetDevice(g_currentDevice));
HIPCHECK(hipMalloc(&A_d0, Nbytes));
HIPCHECK(hipMemset(A_d0, memsetval, Nbytes));
// hipDeviceSynchronize as hipMemset is asynchronous when destination memory is device memory
HIPCHECK(hipDeviceSynchronize());
// allocate and initialize memory on peer device
HIPCHECK(hipSetDevice(g_peerDevice));
HIPCHECK(hipMalloc(&A_d1, Nbytes));
HIPCHECK(hipMemset(A_d1, 0x13, Nbytes));
// hipDeviceSynchronize as hipMemset is asynchronous when destination memory is device memory
HIPCHECK(hipDeviceSynchronize());
//---
@@ -305,6 +312,8 @@ void testPeerHostToDevice(bool useAsyncCopy) {
HIPCHECK(hipSetDevice(g_peerDevice));
HIPCHECK(hipMalloc(&A_d1, Nbytes));
HIPCHECK(hipMemset(A_d1, 0x13, Nbytes));
// hipDeviceSynchronize as hipMemset is asynchronous when destination memory is device memory
HIPCHECK(hipDeviceSynchronize());
bool firstAsyncCopy = useAsyncCopy; /*TODO - should be useAsyncCopy*/