SWDEV-299127 - Merge 'develop' into 'amd-staging'

Change-Id: I486d92f6446aef139235786349d91c15ca816319

.gitignore

@@ -11,6 +11,11 @@ bin/hipBusBandwidth
 bin/hipDispatchLatency
 bin/hipify-clang
 tags
+samples/0_Intro/module_api/runKernel.hip.out
+samples/0_Intro/module_api/vcpy_isa.code
+samples/0_Intro/module_api/vcpy_isa.hsaco
+samples/0_Intro/module_api/vcpy_kernel.co
+samples/0_Intro/module_api/vcpy_kernel.code
 samples/1_Utils/hipInfo/hipInfo
 samples/1_Utils/hipBusBandwidth/hipBusBandwidth
-samples/1_Utils/hipDispatchLatency/hipDispatchLatency
+samples/1_Utils/hipDispatchLatency/hipDispatchLatency

bin/hipcc

@@ -28,7 +28,9 @@ use File::Basename;
 use File::Spec::Functions 'catfile';
 
 # TODO: By default select perl script until change incorporated in HIP build script.
-my $HIPCC_USE_PERL_SCRIPT = 1;
+my $USE_PERL_SCRIPT = $ENV{'HIP_USE_PERL_SCRIPTS'};
+$USE_PERL_SCRIPT //= 1;  # use defined-or assignment operator.  Use env var, but if not defined default to 1.
+
 my $isWindows =  ($^O eq 'MSWin32' or $^O eq 'msys');
 # escapes args with quotes SWDEV-341955
 foreach $arg (@ARGV) {
@@ -38,8 +40,7 @@ foreach $arg (@ARGV) {
 }
 
 my $SCRIPT_DIR=dirname(__FILE__);
-
-if ($HIPCC_USE_PERL_SCRIPT) {
+if ($USE_PERL_SCRIPT) {
     #Invoke hipcc.pl
     my $HIPCC_PERL=catfile($SCRIPT_DIR, '/hipcc.pl');
     system($^X, $HIPCC_PERL, @ARGV);

bin/hipconfig

@@ -28,11 +28,11 @@ use File::Basename;
 use File::Spec::Functions 'catfile';
 
 #TODO: By default select perl script until change incorporated in HIP build script
-my $HIPCONFIG_USE_PERL_SCRIPT = 1;
+my $USE_PERL_SCRIPT = $ENV{'HIP_USE_PERL_SCRIPTS'};
+$USE_PERL_SCRIPT //= 1;  # use defined-or assignment operator.  Use env var, but if not defined default to 1.
 my $isWindows =  ($^O eq 'MSWin32' or $^O eq 'msys');
 my $SCRIPT_DIR=dirname(__FILE__);
-
-if ($HIPCONFIG_USE_PERL_SCRIPT) {
+if ($USE_PERL_SCRIPT) {
     #Invoke hipconfig.pl
     my $HIPCONFIG_PERL=catfile($SCRIPT_DIR, '/hipconfig.pl');
     system($^X, $HIPCONFIG_PERL, @ARGV);

include/hip/hip_runtime_api.h

@@ -6351,7 +6351,7 @@ hipError_t hipUserObjectCreate(hipUserObject_t* object_out, void* ptr, hipHostFn
  * @warning : This API is marked as beta, meaning, while this is feature complete,
  * it is still open to changes and may have outstanding issues.
  */
-hipError_t hipUserObjectRelease(hipUserObject_t object, unsigned int count);
+hipError_t hipUserObjectRelease(hipUserObject_t object, unsigned int count __dparm(1));
 
 /**
  * @brief Retain number of references to resource.
@@ -6362,7 +6362,7 @@ hipError_t hipUserObjectRelease(hipUserObject_t object, unsigned int count);
  * @warning : This API is marked as beta, meaning, while this is feature complete,
  * it is still open to changes and may have outstanding issues.
  */
-hipError_t hipUserObjectRetain(hipUserObject_t object, unsigned int count);
+hipError_t hipUserObjectRetain(hipUserObject_t object, unsigned int count __dparm(1));
 
 /**
  * @brief Retain user object for graphs.
@@ -6375,7 +6375,7 @@ hipError_t hipUserObjectRetain(hipUserObject_t object, unsigned int count);
  * @warning : This API is marked as beta, meaning, while this is feature complete,
  * it is still open to changes and may have outstanding issues.
  */
-hipError_t hipGraphRetainUserObject(hipGraph_t graph, hipUserObject_t object, unsigned int count, unsigned int flags);
+hipError_t hipGraphRetainUserObject(hipGraph_t graph, hipUserObject_t object, unsigned int count __dparm(1), unsigned int flags __dparm(0));
 
 /**
  * @brief Release user object from graphs.

samples/0_Intro/module_api/.gitignore

@@ -1,5 +0,0 @@
-runKernel.hip.out
-vcpy_isa.code
-vcpy_isa.hsaco
-vcpy_kernel.co
-vcpy_kernel.code

tests/catch/include/threaded_zig_zag_test.hh

@@ -0,0 +1,110 @@
+/*
+Copyright (c) 2022 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 WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#pragma once
+
+#include <condition_variable>
+#include <mutex>
+#include <thread>
+
+/*
+Guarantees total ordering between parent and child thread
+PARENT      CHILD
+THREAD      THREAD
+TestPart1
+         \
+          \
+           \
+            TestPart2
+           /
+          /
+         /
+TestPart3
+         \
+          \
+           \
+            TestPart4
+Usage:
+Define a derived class which inherits from ThreadedZigZagTest instantiated with that selfsame class,
+which implements the appropriate test methods
+class DerivedTestClass : public ThreadedZigZagTest<DerivedTestClass> {
+  void TestPart1() {...}
+  void TestPart2() {...}
+  void TestPart3() {...}
+  void TestPart4() {...}
+};
+The derived class can contain state that the test requires.
+*/
+
+template <typename T> class ThreadedZigZagTest {
+ public:
+  void run() {
+    // 1.
+    static_cast<T*>(this)->TestPart1();
+
+    auto t = std::thread([this] {
+      // 2.
+      static_cast<T*>(this)->TestPart2();
+
+      {
+        std::lock_guard<std::mutex> lock(mtx_);
+        ready_ = true;
+      }
+      cv_.notify_one();
+
+      {
+        std::unique_lock<std::mutex> lock(mtx_);
+        cv_.wait(lock, [this] { return !ready_; });
+      }
+
+      // 4.
+      static_cast<T*>(this)->TestPart4();
+    });
+
+    {
+      std::unique_lock<std::mutex> lock(mtx_);
+      cv_.wait(lock, [this] { return ready_; });
+    }
+
+    // 3.
+    static_cast<T*>(this)->TestPart3();
+
+    {
+      std::lock_guard<std::mutex> lock(mtx_);
+      ready_ = false;
+    }
+    cv_.notify_one();
+
+    // Finalize
+    t.join();
+    HIP_CHECK_THREAD_FINALIZE();
+  }
+
+  void TestPart1() const {}
+  void TestPart2() const {}
+  void TestPart3() const {}
+  void TestPart4() const {}
+
+ private:
+  std::mutex mtx_;
+  std::condition_variable cv_;
+  bool ready_ = false;
+};

tests/catch/multiproc/hipIpcMemAccessTest.cc

@@ -200,7 +200,7 @@ TEST_CASE("Unit_hipIpcMemAccess_ParameterValidation") {
   SECTION("Open mem handle with handle as un-initialized") {
     ret = hipIpcOpenMemHandle(&Ad2, MemHandleUninit,
                                            hipIpcMemLazyEnablePeerAccess);
-    REQUIRE(ret == hipErrorInvalidValue);
+    REQUIRE((ret == hipErrorInvalidValue || ret == hipErrorInvalidDevicePointer));
   }
 #if HT_AMD
   // Test is disabled for nvidia as api not returning expected value.

tests/catch/unit/device/CMakeLists.txt

@@ -17,6 +17,7 @@ set(TEST_SRC
     hipSetGetDevice.cc
     hipDeviceGetUuid.cc
     hipDeviceGetP2PAttribute.cc
+    hipExtGetLinkTypeAndHopCount.cc
 )
 
 set_source_files_properties(hipGetDeviceCount.cc PROPERTIES COMPILE_FLAGS -std=c++17)

tests/catch/unit/device/hipExtGetLinkTypeAndHopCount.cc

@@ -0,0 +1,100 @@
+/*
+Copyright (c) 2022 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 WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include <hip_test_common.hh>
+#include <hip/hip_runtime_api.h>
+
+#if HT_AMD
+TEST_CASE("Unit_hipExtGetLinkTypeAndHopCount_Positive_Basic") {
+  const auto device1 = GENERATE(range(0, HipTest::getDeviceCount()));
+  const auto device2 = GENERATE(range(0, HipTest::getDeviceCount()));
+
+  if (device1 == device2) {
+    return;
+  }
+
+  uint32_t link_type1 = -1, hop_count1 = -1;
+  uint32_t link_type2 = -1, hop_count2 = -1;
+
+  HIP_CHECK(hipExtGetLinkTypeAndHopCount(device1, device2, &link_type1, &hop_count1));
+  HIP_CHECK(hipExtGetLinkTypeAndHopCount(device2, device1, &link_type2, &hop_count2));
+
+  REQUIRE(hop_count1 > 0);
+
+  REQUIRE(hop_count1 == hop_count2);
+  REQUIRE(link_type1 == link_type2);
+}
+
+TEST_CASE("Unit_hipExtGetLinkTypeAndHopCount_Negative_Parameters") {
+  uint32_t link_type, hop_count;
+  SECTION("same device") {
+    HIP_CHECK_ERROR(hipExtGetLinkTypeAndHopCount(0, 0, &link_type, &hop_count),
+                    hipErrorInvalidValue);
+  }
+
+  SECTION("device ordinance 1 too large") {
+    HIP_CHECK_ERROR(
+        hipExtGetLinkTypeAndHopCount(HipTest::getDeviceCount(), 0, &link_type, &hop_count),
+        hipErrorInvalidDevice);
+  }
+
+  SECTION("device ordinance 2 too large") {
+    HIP_CHECK_ERROR(
+        hipExtGetLinkTypeAndHopCount(0, HipTest::getDeviceCount(), &link_type, &hop_count),
+        hipErrorInvalidDevice);
+  }
+
+  SECTION("device ordinances too large") {
+    HIP_CHECK_ERROR(
+        hipExtGetLinkTypeAndHopCount(HipTest::getDeviceCount(), HipTest::getDeviceCount() + 1,
+                                     &link_type, &hop_count),
+        hipErrorInvalidDevice);
+  }
+
+  SECTION("device 1 < 0") {
+    HIP_CHECK_ERROR(hipExtGetLinkTypeAndHopCount(-1, 0, &link_type, &hop_count),
+                    hipErrorInvalidValue);
+  }
+
+  SECTION("device 2 < 0") {
+    HIP_CHECK_ERROR(hipExtGetLinkTypeAndHopCount(0, -1, &link_type, &hop_count),
+                    hipErrorInvalidValue);
+  }
+
+  SECTION("both devices < 0") {
+    HIP_CHECK_ERROR(hipExtGetLinkTypeAndHopCount(-1, -2, &link_type, &hop_count),
+                    hipErrorInvalidValue);
+  }
+
+  SECTION("linktype == nullptr") {
+    HIP_CHECK_ERROR(hipExtGetLinkTypeAndHopCount(0, 1, nullptr, &hop_count), hipErrorInvalidValue);
+  }
+
+  SECTION("hopcount == nullptr") {
+    HIP_CHECK_ERROR(hipExtGetLinkTypeAndHopCount(0, 1, &link_type, nullptr), hipErrorInvalidValue);
+  }
+
+  SECTION("linktype and hopcount == nullptr") {
+    HIP_CHECK_ERROR(hipExtGetLinkTypeAndHopCount(0, 1, nullptr, nullptr), hipErrorInvalidValue);
+  }
+}
+#endif

tests/catch/unit/device/hipSetGetDevice.cc

@@ -1,30 +1,34 @@
 /*
- * Copyright (c) 2021 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 WARRANTY OF ANY KIND, EXPRESS OR
- * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
- * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
- * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
- * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
- * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
- * THE SOFTWARE.
- */
+Copyright (c) 2022 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 WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
 
 /*
  * Verifies functionality of hipSetDevice/hipGetDevice api.
  * -- Basic Test to set and get valid device numbers.
  */
 
-#include <hip_test_common.hh>
 #include <thread>
 
+#include <hip_test_common.hh>
+#include <threaded_zig_zag_test.hh>
+
 TEST_CASE("Unit_hipSetDevice_BasicSetGet") {
   int numDevices = 0;
   int device{};
@@ -83,6 +87,44 @@ TEST_CASE("Unit_hipGetSetDevice_MultiThreaded") {
   HIP_CHECK_THREAD_FINALIZE();
 }
 
+TEST_CASE("Unit_hipSetGetDevice_Positive_Threaded_Basic") {
+  class HipSetGetDeviceThreadedTest : public ThreadedZigZagTest<HipSetGetDeviceThreadedTest> {
+   public:
+    void TestPart1() { HIP_CHECK(hipSetDevice(0)); }
+    void TestPart2() {
+      HIP_CHECK_THREAD(hipSetDevice(1));
+      HIP_CHECK_THREAD(hipMalloc(&ptr, 2 * 1024 * 1024));
+    }
+    void TestPart3() {
+      int device = -1;
+      HIP_CHECK_THREAD(hipGetDevice(&device));
+      REQUIRE_THREAD(device == 0);
+      device = -1;
+      // To check if set device worked properly, outside of hipGetDevice
+      HIP_CHECK_THREAD(hipPointerGetAttribute(&device, HIP_POINTER_ATTRIBUTE_DEVICE_ORDINAL,
+                                              reinterpret_cast<hipDeviceptr_t>(ptr)));
+      REQUIRE_THREAD(device == 1);
+    }
+    void TestPart4() {
+      int device = -1;
+      HIP_CHECK_THREAD(hipGetDevice(&device));
+      REQUIRE_THREAD(device == 1);
+      HIP_CHECK_THREAD(hipFree(ptr));
+    }
+
+   private:
+    void* ptr = nullptr;
+  };
+
+  if (HipTest::getDeviceCount() < 2) {
+    HipTest::HIP_SKIP_TEST("This rest requires 2 GPUs. Skipping test");
+    return;
+  }
+
+  HipSetGetDeviceThreadedTest test;
+  test.run();
+}
+
 TEST_CASE("Unit_hipSetGetDevice_Negative") {
   SECTION("Get Device - nullptr") { HIP_CHECK_ERROR(hipGetDevice(nullptr), hipErrorInvalidValue); }
 

tests/catch/unit/memory/hipArray3DCreate.cc

@@ -87,11 +87,6 @@ TEMPLATE_TEST_CASE("Unit_hipArray3DCreate_happy", "", char, uchar2, uint2, int4,
 
 TEMPLATE_TEST_CASE("Unit_hipArray3DCreate_MaxTexture", "", int, uint4, short, ushort2,
                    unsigned char, float, float4) {
-#if HT_AMD
-  HipTest::HIP_SKIP_TEST("EXSWCPHIPT-97");
-  return;
-#endif
-
   using vec_info = vector_info<TestType>;
   DriverContext ctx;
 

tests/catch/unit/streamperthread/hipStreamPerThrdTsts.cc

@@ -155,9 +155,7 @@ __global__ void test_gwsPerThrd(uint* buf, uint bufSize, int64_t* tmpBuf,
     }
 }
 #endif
-static const uint BufferSizeInDwords = 256 * 1024 * 1024;
-static constexpr uint NumKernelArgs = 4;
-static constexpr uint MaxGPUs = 8;
+
 // callback function
 static void HIPRT_CB CallBackFunctn(hipStream_t strm, hipError_t err,
                                     void *ChkVal) {
@@ -472,126 +470,3 @@ TEST_CASE("Unit_hipStreamPerThread_CoopLaunch") {
     }
   }
 }
-
-/* Testing hipLaunchCooperativeKernelMultiDevice() with hipStreamPerThread*/
-#if HT_AMD
-TEST_CASE("Unit_hipStreamPerThread_CoopLaunchMDev") {
-  uint* dA[MaxGPUs];
-  int64_t* dB[MaxGPUs];
-  int64_t* dC;
-
-  uint32_t* init = new uint32_t[BufferSizeInDwords];
-  for (uint32_t i = 0; i < BufferSizeInDwords; ++i) {
-    init[i] = i;
-  }
-
-  int nGpu = 0;
-  HIPCHECK(hipGetDeviceCount(&nGpu));
-  size_t copySizeInDwords = BufferSizeInDwords / nGpu;
-  hipDeviceProp_t deviceProp[MaxGPUs];
-
-  for (int i = 0; i < nGpu; i++) {
-    HIPCHECK(hipSetDevice(i));
-
-    // Calculate the device occupancy to know how many blocks can be
-    // run concurrently
-    HIP_CHECK(hipGetDeviceProperties(&deviceProp[i], 0));
-    if (!deviceProp[i].cooperativeMultiDeviceLaunch) {
-      WARN("Device doesn't support cooperative launch!");
-      return;
-    }
-    size_t SIZE = copySizeInDwords * sizeof(uint);
-
-    HIPCHECK(hipMalloc(reinterpret_cast<void**>(&dA[i]), SIZE));
-    HIPCHECK(hipMalloc(reinterpret_cast<void**>(&dB[i]),
-             64 * deviceProp[i].multiProcessorCount * sizeof(int64_t)));
-    if (i == 0) {
-      HIPCHECK(hipHostMalloc(reinterpret_cast<void**>(&dC),
-      (nGpu + 1) * sizeof(int64_t)));
-    }
-    HIPCHECK(hipMemcpy(dA[i], &init[i * copySizeInDwords] , SIZE,
-             hipMemcpyHostToDevice));
-    HIP_CHECK(hipDeviceSynchronize());
-  }
-
-  dim3 dimBlock;
-  dim3 dimGrid;
-  dimGrid.x = 1;
-  dimGrid.y = 1;
-  dimGrid.z = 1;
-  dimBlock.x = 64;
-  dimBlock.y = 1;
-  dimBlock.z = 1;
-
-  int numBlocks = 0;
-  uint workgroups[3] = {64, 128, 256};
-
-  hipLaunchParams* launchParamsList = new hipLaunchParams[nGpu];
-  std::time_t end_time;
-  double time = 0;
-  for (uint set = 0; set < 3; ++set) {
-    void* args[MaxGPUs * NumKernelArgs];
-    WARN("---------- Test#" << set << ", size: "<< BufferSizeInDwords <<
-      " dwords ---------------\n");
-    for (int i = 0; i < nGpu; i++) {
-      HIPCHECK(hipSetDevice(i));
-      dimBlock.x = workgroups[set];
-      HIPCHECK(hipOccupancyMaxActiveBlocksPerMultiprocessor(&numBlocks,
-      test_gwsPerThrd, dimBlock.x * dimBlock.y * dimBlock.z,
-      dimBlock.x * sizeof(int64_t)));
-
-      WARN("GPU(" << i << ") Block size: " << dimBlock.x <<
-        " Num blocks per CU: " << numBlocks << "\n");
-
-      dimGrid.x = deviceProp[i].multiProcessorCount * (std::min)(numBlocks, 32);
-
-      args[i * NumKernelArgs]     = reinterpret_cast<void*>(&dA[i]);
-      args[i * NumKernelArgs + 1] = reinterpret_cast<void*>(&copySizeInDwords);
-      args[i * NumKernelArgs + 2] = reinterpret_cast<void*>(&dB[i]);
-      args[i * NumKernelArgs + 3] = reinterpret_cast<void*>(&dC);
-
-      launchParamsList[i].func = reinterpret_cast<void*>(test_gwsPerThrd);
-      launchParamsList[i].gridDim = dimGrid;
-      launchParamsList[i].blockDim = dimBlock;
-      launchParamsList[i].sharedMem = dimBlock.x * sizeof(int64_t);
-
-      launchParamsList[i].stream = hipStreamPerThread;
-      launchParamsList[i].args = &args[i * NumKernelArgs];
-    }
-
-    system_clock::time_point start = system_clock::now();
-    HIP_CHECK(hipLaunchCooperativeKernelMultiDevice(launchParamsList, nGpu, 0));
-    for (int i = 0; i < nGpu; i++) {
-      HIP_CHECK(hipSetDevice(i));
-      HIP_CHECK(hipDeviceSynchronize());
-    }
-    system_clock::time_point end = system_clock::now();
-    std::chrono::duration<double> elapsed_seconds = end - start;
-    end_time = std::chrono::system_clock::to_time_t(end);
-
-    time += elapsed_seconds.count();
-
-    size_t processedDwords = copySizeInDwords * nGpu;
-    if (*dC != (((int64_t)(processedDwords) * (processedDwords - 1)) / 2)) {
-      WARN("Data validation failed ("<< *dC << " != " <<
-        (((int64_t)(BufferSizeInDwords) * (BufferSizeInDwords - 1)) / 2) <<
-        ") for grid size = " << dimGrid.x << " and block size = " <<
-        dimBlock.x << "\n");
-      WARN("Test failed!");
-    }
-  }
-
-  delete [] launchParamsList;
-
-  WARN("finished computation at " << std::ctime(&end_time));
-  WARN("elapsed time: " << time << "s\n");
-
-  HIP_CHECK(hipSetDevice(0));
-  HIP_CHECK(hipFree(dC));
-  for (int i = 0; i < nGpu; i++) {
-    HIP_CHECK(hipFree(dA[i]));
-    HIP_CHECK(hipFree(dB[i]));
-  }
-  delete [] init;
-}
-#endif

tests/src/p2p/hipPeerToPeer_simple.cpp

@@ -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*/