diff --git a/.gitignore b/.gitignore index fe07943cad..6bdb3a4030 100644 --- a/.gitignore +++ b/.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 \ No newline at end of file +samples/1_Utils/hipDispatchLatency/hipDispatchLatency diff --git a/bin/hipcc b/bin/hipcc index 62191ffa4e..42fa1e8a20 100755 --- a/bin/hipcc +++ b/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); diff --git a/bin/hipconfig b/bin/hipconfig index 6bb4f892f8..d50b10b94a 100755 --- a/bin/hipconfig +++ b/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); diff --git a/include/hip/hip_runtime_api.h b/include/hip/hip_runtime_api.h index 86bef4c273..6836fa5358 100644 --- a/include/hip/hip_runtime_api.h +++ b/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. diff --git a/samples/0_Intro/module_api/.gitignore b/samples/0_Intro/module_api/.gitignore deleted file mode 100644 index c1d81e043f..0000000000 --- a/samples/0_Intro/module_api/.gitignore +++ /dev/null @@ -1,5 +0,0 @@ -runKernel.hip.out -vcpy_isa.code -vcpy_isa.hsaco -vcpy_kernel.co -vcpy_kernel.code diff --git a/tests/catch/include/threaded_zig_zag_test.hh b/tests/catch/include/threaded_zig_zag_test.hh new file mode 100644 index 0000000000..96fe250869 --- /dev/null +++ b/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 +#include +#include + +/* +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 { + void TestPart1() {...} + void TestPart2() {...} + void TestPart3() {...} + void TestPart4() {...} +}; +The derived class can contain state that the test requires. +*/ + +template class ThreadedZigZagTest { + public: + void run() { + // 1. + static_cast(this)->TestPart1(); + + auto t = std::thread([this] { + // 2. + static_cast(this)->TestPart2(); + + { + std::lock_guard lock(mtx_); + ready_ = true; + } + cv_.notify_one(); + + { + std::unique_lock lock(mtx_); + cv_.wait(lock, [this] { return !ready_; }); + } + + // 4. + static_cast(this)->TestPart4(); + }); + + { + std::unique_lock lock(mtx_); + cv_.wait(lock, [this] { return ready_; }); + } + + // 3. + static_cast(this)->TestPart3(); + + { + std::lock_guard 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; +}; \ No newline at end of file diff --git a/tests/catch/multiproc/hipIpcMemAccessTest.cc b/tests/catch/multiproc/hipIpcMemAccessTest.cc index 6dd0941b4a..fee23c1ef4 100644 --- a/tests/catch/multiproc/hipIpcMemAccessTest.cc +++ b/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. diff --git a/tests/catch/unit/device/CMakeLists.txt b/tests/catch/unit/device/CMakeLists.txt index 7d72853100..b0b1a495b9 100644 --- a/tests/catch/unit/device/CMakeLists.txt +++ b/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) diff --git a/tests/catch/unit/device/hipExtGetLinkTypeAndHopCount.cc b/tests/catch/unit/device/hipExtGetLinkTypeAndHopCount.cc new file mode 100644 index 0000000000..11b0faa98f --- /dev/null +++ b/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 +#include + +#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 \ No newline at end of file diff --git a/tests/catch/unit/device/hipSetGetDevice.cc b/tests/catch/unit/device/hipSetGetDevice.cc index 011779a489..62df5f5ac8 100644 --- a/tests/catch/unit/device/hipSetGetDevice.cc +++ b/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 #include +#include +#include + 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 { + 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(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); } diff --git a/tests/catch/unit/memory/hipArray3DCreate.cc b/tests/catch/unit/memory/hipArray3DCreate.cc index 9c11152936..46c31c1b52 100644 --- a/tests/catch/unit/memory/hipArray3DCreate.cc +++ b/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; DriverContext ctx; diff --git a/tests/catch/unit/streamperthread/hipStreamPerThrdTsts.cc b/tests/catch/unit/streamperthread/hipStreamPerThrdTsts.cc index 418e5332ab..95f5c27fd5 100644 --- a/tests/catch/unit/streamperthread/hipStreamPerThrdTsts.cc +++ b/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(&dA[i]), SIZE)); - HIPCHECK(hipMalloc(reinterpret_cast(&dB[i]), - 64 * deviceProp[i].multiProcessorCount * sizeof(int64_t))); - if (i == 0) { - HIPCHECK(hipHostMalloc(reinterpret_cast(&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(&dA[i]); - args[i * NumKernelArgs + 1] = reinterpret_cast(©SizeInDwords); - args[i * NumKernelArgs + 2] = reinterpret_cast(&dB[i]); - args[i * NumKernelArgs + 3] = reinterpret_cast(&dC); - - launchParamsList[i].func = reinterpret_cast(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 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 diff --git a/tests/src/p2p/hipPeerToPeer_simple.cpp b/tests/src/p2p/hipPeerToPeer_simple.cpp index f2f3592df2..1f86fe4eb5 100644 --- a/tests/src/p2p/hipPeerToPeer_simple.cpp +++ b/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*/