From 9d9c36e829542a559ed54ad7882ae59d5851ab56 Mon Sep 17 00:00:00 2001 From: Jatin Chaudhary Date: Wed, 3 Aug 2022 12:17:34 +0100 Subject: [PATCH 1/9] Add tests for hipStreamWaitEvent (#2559) --- tests/catch/CMakeLists.txt | 5 + tests/catch/unit/stream/CMakeLists.txt | 3 +- tests/catch/unit/stream/hipStreamWaitEvent.cc | 121 ++++++++++++++++++ 3 files changed, 128 insertions(+), 1 deletion(-) create mode 100644 tests/catch/unit/stream/hipStreamWaitEvent.cc diff --git a/tests/catch/CMakeLists.txt b/tests/catch/CMakeLists.txt index a8a488002d..f9e11b9c82 100644 --- a/tests/catch/CMakeLists.txt +++ b/tests/catch/CMakeLists.txt @@ -201,6 +201,11 @@ set(_catchInfo ${_catchInfo} "HIP_PLATFORM=${HIP_PLATFORM}\n") set(_catchInfo ${_catchInfo} "ARCHS=${HIP_GPU_ARCH_LIST}\n") file(WRITE ${CMAKE_CURRENT_BINARY_DIR}/${CATCH_BUILD_DIR}/catchInfo.txt ${_catchInfo}) +# Enable device lambda on nvidia platforms +if(HIP_COMPILER MATCHES "nvcc") + set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} --extended-lambda") +endif() + # Disable CXX extensions (gnu++11 etc) set(CMAKE_CXX_EXTENSIONS OFF) diff --git a/tests/catch/unit/stream/CMakeLists.txt b/tests/catch/unit/stream/CMakeLists.txt index 229df157c1..6cdbac9a4e 100644 --- a/tests/catch/unit/stream/CMakeLists.txt +++ b/tests/catch/unit/stream/CMakeLists.txt @@ -13,6 +13,7 @@ set(TEST_SRC hipStreamValue.cc hipStreamWithCUMask.cc hipStreamACb_MultiThread.cc + hipStreamWaitEvent.cc ) else() set(TEST_SRC @@ -31,7 +32,7 @@ set(TEST_SRC hipStreamValue.cc ) - set_source_files_properties(hipStreamAttachMemAsync.cc PROPERTIES COMPILE_FLAGS -std=c++17) + # set_source_files_properties(hipStreamAttachMemAsync.cc PROPERTIES COMPILE_FLAGS -std=c++17) endif() hip_add_exe_to_target(NAME StreamTest diff --git a/tests/catch/unit/stream/hipStreamWaitEvent.cc b/tests/catch/unit/stream/hipStreamWaitEvent.cc new file mode 100644 index 0000000000..cea39fcd54 --- /dev/null +++ b/tests/catch/unit/stream/hipStreamWaitEvent.cc @@ -0,0 +1,121 @@ +/* +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 + +TEST_CASE("Unit_hipStreamWaitEvent_Negative") { + enum class StreamTestType { NullStream = 0, StreamPerThread, CreatedStream }; + + auto streamType = GENERATE(StreamTestType::NullStream, StreamTestType::StreamPerThread, + StreamTestType::CreatedStream); + + hipStream_t stream{nullptr}; + hipEvent_t event{nullptr}; + + if (streamType == StreamTestType::StreamPerThread) { + stream = hipStreamPerThread; + } else if (streamType == StreamTestType::CreatedStream) { + HIP_CHECK(hipStreamCreate(&stream)); + } + + HIP_CHECK(hipEventCreate(&event)); + + REQUIRE((stream != nullptr) != (streamType == StreamTestType::NullStream)); + REQUIRE(event != nullptr); + + SECTION("Invalid Event") { + INFO("Running against Invalid Event"); + HIP_CHECK_ERROR(hipStreamWaitEvent(stream, nullptr, 0), hipErrorInvalidResourceHandle); + } + + SECTION("Invalid Flags") { + INFO("Running against Invalid Flags"); + constexpr unsigned flag = ~0u; + REQUIRE(flag != 0); + HIP_CHECK_ERROR(hipStreamWaitEvent(stream, event, flag), hipErrorInvalidValue); + } + + HIP_CHECK(hipEventDestroy(event)); + + if (streamType == StreamTestType::CreatedStream) { + HIP_CHECK(hipStreamDestroy(stream)); + } +} + +// Since we can not use atomic*_system on every gpu, we will use wait based on clock rate. +// This wont be accurate since clock rate of a GPU varies depending on many variables including +// thermals, load, utilization +__global__ void waitKernel(int clockRate, int seconds) { + auto start = clock(); + auto ms = seconds * 1000; + long long waitTill = clockRate * (long long)ms; + while (1) { + auto end = clock(); + if ((end - start) > waitTill) { + return; + } + } +} + +TEST_CASE("Unit_hipStreamWaitEvent_DifferentStreams") { + hipStream_t blockedStreamA{nullptr}, streamBlockedOnStreamA{nullptr}, unblockingStream{nullptr}; + hipEvent_t waitEvent{nullptr}; + + HIP_CHECK(hipStreamCreate(&blockedStreamA)); + HIP_CHECK(hipStreamCreate(&streamBlockedOnStreamA)); + HIP_CHECK(hipStreamCreate(&unblockingStream)); + HIP_CHECK(hipEventCreate(&waitEvent)); + + REQUIRE(blockedStreamA != nullptr); + REQUIRE(streamBlockedOnStreamA != nullptr); + REQUIRE(waitEvent != nullptr); + + hipDeviceProp_t prop{}; + HIP_CHECK(hipGetDeviceProperties(&prop, 0)); + auto clockRate = prop.clockRate; + + waitKernel<<<1, 1, 0, blockedStreamA>>>(clockRate, + 3); // wait for 5 seconds + HIP_CHECK(hipEventRecord(waitEvent, blockedStreamA)); + + // Make sure stream is waiting for data to be set + HIP_CHECK_ERROR(hipEventQuery(waitEvent), hipErrorNotReady); + + HIP_CHECK(hipStreamWaitEvent(streamBlockedOnStreamA, waitEvent, 0)); + + waitKernel<<<1, 1, 0, streamBlockedOnStreamA>>>(clockRate, 2); // Wait for 2 seconds + + // Make sure stream is waiting for event on blockedStreamA + HIP_CHECK_ERROR(hipStreamQuery(streamBlockedOnStreamA), hipErrorNotReady); + + HIP_CHECK(hipStreamSynchronize(unblockingStream)); + + HIP_CHECK(hipStreamSynchronize(blockedStreamA)); + HIP_CHECK(hipStreamSynchronize(streamBlockedOnStreamA)); + + // Check that both streams have finished + HIP_CHECK(hipStreamQuery(blockedStreamA)); + HIP_CHECK(hipStreamQuery(streamBlockedOnStreamA)); + + HIP_CHECK(hipStreamDestroy(blockedStreamA)); + HIP_CHECK(hipStreamDestroy(streamBlockedOnStreamA)); + HIP_CHECK(hipStreamDestroy(unblockingStream)); + HIP_CHECK(hipEventDestroy(waitEvent)); +} From c435fd5c473691f3116f7612ae4ffd3f4957aae7 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?F=C3=A1bio?= Date: Wed, 3 Aug 2022 12:19:49 +0100 Subject: [PATCH 2/9] EXSWCPHIPT-123: Added testing for hipDeviceGetP2PAttribute (#2840) --- tests/catch/unit/device/CMakeLists.txt | 4 + .../unit/device/hipDeviceGetP2PAttribute.cc | 126 ++++++++++++++++++ .../device/hipDeviceGetP2PAttribute_exe.cc | 66 +++++++++ 3 files changed, 196 insertions(+) create mode 100644 tests/catch/unit/device/hipDeviceGetP2PAttribute.cc create mode 100644 tests/catch/unit/device/hipDeviceGetP2PAttribute_exe.cc diff --git a/tests/catch/unit/device/CMakeLists.txt b/tests/catch/unit/device/CMakeLists.txt index 247867bf5c..7d72853100 100644 --- a/tests/catch/unit/device/CMakeLists.txt +++ b/tests/catch/unit/device/CMakeLists.txt @@ -16,11 +16,14 @@ set(TEST_SRC hipGetSetDeviceFlags.cc hipSetGetDevice.cc hipDeviceGetUuid.cc + hipDeviceGetP2PAttribute.cc ) set_source_files_properties(hipGetDeviceCount.cc PROPERTIES COMPILE_FLAGS -std=c++17) +set_source_files_properties(hipDeviceGetP2PAttribute.cc PROPERTIES COMPILE_FLAGS -std=c++17) add_executable(getDeviceCount EXCLUDE_FROM_ALL getDeviceCount_exe.cc) +add_executable(hipDeviceGetP2PAttribute EXCLUDE_FROM_ALL hipDeviceGetP2PAttribute_exe.cc) hip_add_exe_to_target(NAME DeviceTest TEST_SRC ${TEST_SRC} @@ -28,3 +31,4 @@ hip_add_exe_to_target(NAME DeviceTest COMPILE_OPTIONS -std=c++14) add_dependencies(DeviceTest getDeviceCount) +add_dependencies(DeviceTest hipDeviceGetP2PAttribute) diff --git a/tests/catch/unit/device/hipDeviceGetP2PAttribute.cc b/tests/catch/unit/device/hipDeviceGetP2PAttribute.cc new file mode 100644 index 0000000000..805eb117ff --- /dev/null +++ b/tests/catch/unit/device/hipDeviceGetP2PAttribute.cc @@ -0,0 +1,126 @@ +/* +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 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. +*/ + +#include +#include +#include +#include "hip/hip_runtime_api.h" +#include +#include + +/** + * @brief Test all possible combination of attributes and devices for hipDeviceGetP2PAttribute. + * Verify that the output is within the range of acceptable values. + * + */ +TEST_CASE("Unit_hipDeviceGetP2PAttribute_Basic") { +#if HT_AMD + HipTest::HIP_SKIP_TEST("EXSWCPHIPT-119"); + return; +#else + + int deviceCount = HipTest::getGeviceCount(); + if (deviceCount < 2) { + HipTest::HIP_SKIP_TEST("Skipping because devices < 2"); + return; + } + + hipDeviceP2PAttr attribute = + GENERATE(hipDevP2PAttrPerformanceRank, hipDevP2PAttrAccessSupported, + hipDevP2PAttrNativeAtomicSupported, hipDevP2PAttrHipArrayAccessSupported); + + /* Test all combinations of devices in the system */ + for (int srcDevice = 0; srcDevice < deviceCount; ++srcDevice) { + for (int dstDevice = 0; dstDevice < deviceCount; ++dstDevice) { + if (srcDevice != dstDevice) { + int value{-1}; + HIP_CHECK(hipDeviceGetP2PAttribute(&value, attribute, srcDevice, dstDevice)); + INFO("hipDeviceP2PAttr: " << attribute << "\nsrcDevice: " << srcDevice + << "\ndstDevice: " << dstDevice << "\nValue: " << value); + if (attribute == hipDevP2PAttrPerformanceRank) { + REQUIRE(value >= 0); + } else { + REQUIRE((value == 0 || value == 1)); + } + } + } + } +#endif +} + +/** + * @brief Negative test scenarios for hipDeviceGetP2PAttribute + * + */ +TEST_CASE("Unit_hipDeviceGetP2PAttribute_Negative") { +#if HT_AMD + HipTest::HIP_SKIP_TEST("EXSWCPHIPT-122"); + return; +#else + + int deviceCount = HipTest::getGeviceCount(); + if (deviceCount < 2) { + HipTest::HIP_SKIP_TEST("Skipping because devices < 2"); + return; + } + + int value; + int validSrcDevice = 0; + int validDstDevice = 1; + hipDeviceP2PAttr validAttr = hipDevP2PAttrAccessSupported; + + SECTION("Nullptr value") { + HIP_CHECK_ERROR(hipDeviceGetP2PAttribute(nullptr, validAttr, validSrcDevice, validDstDevice), + hipErrorInvalidValue); + } + + SECTION("Invalid attribute") { + hipDeviceP2PAttr invalidAttr = static_cast(10); + HIP_CHECK_ERROR(hipDeviceGetP2PAttribute(&value, invalidAttr, validSrcDevice, validDstDevice), + hipErrorInvalidValue); + } + + SECTION("Invalid device") { + int invalidDevice = -1; + HIP_CHECK_ERROR(hipDeviceGetP2PAttribute(&value, validAttr, invalidDevice, validDstDevice), + hipErrorInvalidDevice); + HIP_CHECK_ERROR(hipDeviceGetP2PAttribute(&value, validAttr, validSrcDevice, invalidDevice), + hipErrorInvalidDevice); + } + + /* https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#env-vars */ + SECTION("Hidden devices using environment variables") { + REQUIRE(hip::SpawnProc("hipDeviceGetP2PAttribute").run("") == hipSuccess); + REQUIRE(hip::SpawnProc("hipDeviceGetP2PAttribute").run("0") == hipErrorInvalidDevice); + REQUIRE(hip::SpawnProc("hipDeviceGetP2PAttribute").run("1") == hipErrorInvalidDevice); + REQUIRE(hip::SpawnProc("hipDeviceGetP2PAttribute").run("0,1") == hipSuccess); + REQUIRE(hip::SpawnProc("hipDeviceGetP2PAttribute").run("-1,0") == hipErrorNoDevice); + REQUIRE(hip::SpawnProc("hipDeviceGetP2PAttribute").run("0,-1") == hipErrorInvalidDevice); + REQUIRE(hip::SpawnProc("hipDeviceGetP2PAttribute").run("0,1,-1") == hipSuccess); + REQUIRE(hip::SpawnProc("hipDeviceGetP2PAttribute").run("0,-1,1") == hipErrorInvalidDevice); + + if (deviceCount > 2) { + REQUIRE(hip::SpawnProc("hipDeviceGetP2PAttribute").run("2,1") == hipSuccess); + REQUIRE(hip::SpawnProc("hipDeviceGetP2PAttribute").run("2") == hipErrorInvalidDevice); + } else { + REQUIRE(hip::SpawnProc("hipDeviceGetP2PAttribute").run("2,1") == hipErrorNoDevice); + } + } +#endif +} \ No newline at end of file diff --git a/tests/catch/unit/device/hipDeviceGetP2PAttribute_exe.cc b/tests/catch/unit/device/hipDeviceGetP2PAttribute_exe.cc new file mode 100644 index 0000000000..44e98495ca --- /dev/null +++ b/tests/catch/unit/device/hipDeviceGetP2PAttribute_exe.cc @@ -0,0 +1,66 @@ +/* +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 "hip/hip_runtime_api.h" +#include + + +void inline hideDevices(const char* devices) { +#if HT_NVIDIA + setenv("CUDA_VISIBLE_DEVICES", devices, 1); +#else + setenv("HIP_VISIBLE_DEVICES", devices, 1); + setenv("ROCR_VISIBLE_DEVICES", devices, 1); +#endif +} + +void inline unhideAllDevices() { +#if HT_NVIDIA + unsetenv("CUDA_VISIBLE_DEVICES"); +#else + unsetenv("HIP_VISIBLE_DEVICES"); + unsetenv("ROCR_VISIBLE_DEVICES"); +#endif +} + +/** + * @brief Runs hipDeviceGetP2PAttribute with srcDevice = 0 and dstDevice = 1 + * Expects 1 command line arg, which is the Device Visible String + * + * @return the error code returned by hipDeviceGetP2PAttribute + */ +int main(int argc, char** argv) { + int value; + const int srcDevice = 0; + const int dstDevice = 1; + const hipDeviceP2PAttr validAttr = hipDevP2PAttrAccessSupported; + + if (argc == 2) { + hideDevices(argv[1]); + } + + hipError_t error = hipDeviceGetP2PAttribute(&value, validAttr, srcDevice, dstDevice); + + if (argc == 2) { + unhideAllDevices(); + } + + return error; +} \ No newline at end of file From 7b23c605bf0ad00c684813c354296ac484797b91 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?F=C3=A1bio?= Date: Wed, 3 Aug 2022 12:20:07 +0100 Subject: [PATCH 3/9] EXSWCPHIPT-138: Update hipHostRegister flag testing (#2835) --- tests/catch/unit/memory/hipHostRegister.cc | 43 ++++++++++++++++++---- 1 file changed, 36 insertions(+), 7 deletions(-) diff --git a/tests/catch/unit/memory/hipHostRegister.cc b/tests/catch/unit/memory/hipHostRegister.cc index 417ab0857a..3ccc79d834 100644 --- a/tests/catch/unit/memory/hipHostRegister.cc +++ b/tests/catch/unit/memory/hipHostRegister.cc @@ -29,6 +29,7 @@ This testfile verifies the following scenarios of hipHostRegister API #include #include +#include "hip/hip_runtime_api.h" #define OFFSET 128 static constexpr auto LEN{1024 * 1024}; @@ -160,6 +161,41 @@ template __global__ void fill_kernel(T* dataPtr, T value) { dataPtr[tid] = value; } +TEMPLATE_TEST_CASE("Unit_hipHostRegister_Flags", "", int, float, double) { +#if HT_AMD + HipTest::HIP_SKIP_TEST("EXSWCPHIPT-138"); + return; +#else + + size_t sizeBytes = 1 * sizeof(TestType); + TestType* hostPtr = reinterpret_cast(malloc(sizeBytes)); + + /* Flags aren't used for AMD devices currently */ + struct FlagType { + unsigned int value; + bool valid; + }; + + /* EXSWCPHIPT-29 - 0x08 is hipHostRegisterReadOnly which currently doesn't have a definition in the headers */ + /* hipHostRegisterIoMemory is a valid flag but requires access to I/O mapped memory to be tested */ + FlagType flags = GENERATE( + FlagType{hipHostRegisterDefault, true}, FlagType{hipHostRegisterPortable, true}, + FlagType{0x08, true}, FlagType{hipHostRegisterPortable | hipHostRegisterMapped, true}, + FlagType{hipHostRegisterPortable | hipHostRegisterMapped | 0x08, true}, FlagType{0xF0, false}, + FlagType{0xFFF2, false}, FlagType{0xFFFFFFFF, false}); + + INFO("Testing hipHostRegister flag: " << flags.value); + if (flags.valid) { + HIP_CHECK(hipHostRegister(hostPtr, sizeBytes, flags.value)); + HIP_CHECK(hipHostUnregister(hostPtr)); + } else { + HIP_CHECK_ERROR(hipHostRegister(hostPtr, sizeBytes, flags.value), hipErrorInvalidValue); + } + + free(hostPtr); +#endif +} + TEMPLATE_TEST_CASE("Unit_hipHostRegister_Negative", "", int, float, double) { TestType* hostPtr = nullptr; @@ -173,13 +209,6 @@ TEMPLATE_TEST_CASE("Unit_hipHostRegister_Negative", "", int, float, double) { HIP_CHECK_ERROR(hipHostRegister(hostPtr, 0, 0), hipErrorInvalidValue); } -#if HT_NVIDIA - // Flags aren't used for AMD devices currently - SECTION("hipHostRegister Negative Test - invalid flag") { - HIP_CHECK_ERROR(hipHostRegister(hostPtr, sizeBytes, 0b11111111), hipErrorInvalidValue); - } -#endif - size_t devMemAvail{0}, devMemFree{0}; HIP_CHECK(hipMemGetInfo(&devMemFree, &devMemAvail)); auto hostMemFree = HipTest::getMemoryAmount() /* In MB */ * 1024 * 1024; // In bytes From 1a3cbc472e91be4348415fe7947f246543cb653a Mon Sep 17 00:00:00 2001 From: Jatin Chaudhary Date: Wed, 3 Aug 2022 14:54:02 +0100 Subject: [PATCH 4/9] EXSWCPHIPT-35 - Add hipEventQuery tests (#2721) --- tests/catch/unit/event/CMakeLists.txt | 8 ++ tests/catch/unit/event/Unit_hipEventQuery.cc | 97 ++++++++++++++++++++ 2 files changed, 105 insertions(+) create mode 100644 tests/catch/unit/event/Unit_hipEventQuery.cc diff --git a/tests/catch/unit/event/CMakeLists.txt b/tests/catch/unit/event/CMakeLists.txt index fae1b4264c..bbcc2c9fa4 100644 --- a/tests/catch/unit/event/CMakeLists.txt +++ b/tests/catch/unit/event/CMakeLists.txt @@ -7,6 +7,14 @@ set(TEST_SRC Unit_hipEventIpc.cc ) +# The test used wait mechanism and doesnt play well with all arch of nvidia +if(HIP_PLATFORM MATCHES "amd") + set(AMD_SRC + Unit_hipEventQuery.cc + ) + set(TEST_SRC ${TEST_SRC} ${AMD_SRC}) +endif() + hip_add_exe_to_target(NAME EventTest TEST_SRC ${TEST_SRC} TEST_TARGET_NAME build_tests) diff --git a/tests/catch/unit/event/Unit_hipEventQuery.cc b/tests/catch/unit/event/Unit_hipEventQuery.cc new file mode 100644 index 0000000000..d355a515e5 --- /dev/null +++ b/tests/catch/unit/event/Unit_hipEventQuery.cc @@ -0,0 +1,97 @@ +/* +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 + +// Since we can not use atomic*_system on every gpu, we will use wait based on clock rate. +// This wont be accurate since current clock rate of a GPU varies depending on many variables +// including thermals, load, utilization etc +__global__ void waitKernel(int clockRate, int seconds) { + auto start = clock(); + auto ms = seconds * 1000; + long long waitTill = clockRate * (long long)ms; + while (1) { + auto end = clock(); + if ((end - start) > waitTill) { + return; + } + } +} + +TEST_CASE("Unit_hipEventQuery_DifferentDevice") { + hipEvent_t event1{}, event2{}; + HIP_CHECK(hipEventCreate(&event1)); + HIP_CHECK(hipEventCreate(&event2)); + REQUIRE(event1 != nullptr); + REQUIRE(event2 != nullptr); + + hipStream_t stream{nullptr}; + HIP_CHECK(hipStreamCreate(&stream)); + REQUIRE(stream != nullptr); + + hipDeviceProp_t prop{}; + HIP_CHECK(hipGetDeviceProperties(&prop, 0)); + auto clockRate = prop.clockRate; + + // Start kernel on 1st device + { + HIP_CHECK(hipSetDevice(0)); + HIP_CHECK(hipEventRecord(event1, stream)); + + // Start kernel and wait for 3 seconds + // Make sure you increase this time if you add more tests here + waitKernel<<<1, 1, 0, stream>>>(clockRate, 3); + + HIP_CHECK(hipEventRecord(event2, stream)); + + HIP_CHECK(hipEventSynchronize(event1)); + HIP_CHECK(hipEventQuery(event1)); // Should be done + + HIP_CHECK_ERROR(hipEventQuery(event2), + hipErrorNotReady); // Wont be done since kernel is waiting + } + + // If other devices are available, set it + if (HipTest::getDeviceCount() > 1) { + HIP_CHECK(hipSetDevice(1)); + } + + // Query from same or other device depending on availability + { + HIP_CHECK(hipEventQuery(event1)); + HIP_CHECK_ERROR(hipEventQuery(event2), hipErrorNotReady); + + // Sync + // if it hangs here it means GPU kernel cant see the value update by atomic store + HIP_CHECK(hipEventSynchronize(event2)); + + // Query, should be done now + HIP_CHECK(hipEventQuery(event2)); + } + + // Query on same device if multiple devices are present + if (HipTest::getDeviceCount() > 1) { + HIP_CHECK(hipSetDevice(0)); + HIP_CHECK(hipEventQuery(event2)); + } + + HIP_CHECK(hipEventDestroy(event1)); + HIP_CHECK(hipEventDestroy(event2)); + HIP_CHECK(hipStreamDestroy(stream)); +} From 357186bf9e6df79dbdd48570fc52ac892e2a1ec7 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Wed, 3 Aug 2022 23:45:15 +0530 Subject: [PATCH 5/9] SWDEV-346857 - Removed hard coded path /opt/rocm (#2839) As per the requirements in SWDEV-287540 ,config files should not have hard coded path /opt/rocm. Removed the same Also handled the case where ROCM_PATH is defined and not set in ENV. Earlier this condition was getting handled in default case Change-Id: I9b28c4208175be61bc7c90c1fe2885fd299c5536 --- cmake/FindHIP.cmake | 17 ++++++++--------- hip-lang-config.cmake.in | 4 +--- 2 files changed, 9 insertions(+), 12 deletions(-) diff --git a/cmake/FindHIP.cmake b/cmake/FindHIP.cmake index d3d53a9273..8b196dea65 100644 --- a/cmake/FindHIP.cmake +++ b/cmake/FindHIP.cmake @@ -63,7 +63,6 @@ if(NOT APPLE) "$ENV{ROCM_PATH}/hip" ENV HIP_PATH ${_IMPORT_PREFIX} - /opt/rocm/hip DOC "HIP installed location" NO_DEFAULT_PATH ) @@ -86,8 +85,6 @@ if(NOT APPLE) "${HIP_ROOT_DIR}" ENV ROCM_PATH ENV HIP_PATH - /opt/rocm - /opt/rocm/hip PATH_SUFFIXES bin NO_DEFAULT_PATH ) @@ -104,8 +101,6 @@ if(NOT APPLE) "${HIP_ROOT_DIR}" ENV ROCM_PATH ENV HIP_PATH - /opt/rocm - /opt/rocm/hip PATH_SUFFIXES bin NO_DEFAULT_PATH ) @@ -128,8 +123,6 @@ if(NOT APPLE) "${HIP_ROOT_DIR}" ENV ROCM_PATH ENV HIP_PATH - /opt/rocm - /opt/rocm/hip PATH_SUFFIXES bin NO_DEFAULT_PATH ) @@ -253,8 +246,11 @@ elseif("${HIP_COMPILER}" STREQUAL "clang") else() set(HIP_CLANG_PATH "${HIP_PATH}/../llvm/bin") endif() + # Handle the case where ROCM_PATH is defined and not set in ENV + elseif(DEFINED ROCM_PATH) + set(HIP_CLANG_PATH "${ROCM_PATH}/llvm/bin") else() - set(HIP_CLANG_PATH "/opt/rocm/llvm/bin") + message(FATAL_ERROR "Unable to find the clang compiler path. Set ROCM_PATH or HIP_PATH in env ") endif() endif() #Number of parallel jobs by default is 1 @@ -680,8 +676,11 @@ macro(HIP_ADD_EXECUTABLE hip_target) else() set(HIP_CLANG_PATH "${HIP_PATH}/../llvm/bin") endif() + # Handle the case where ROCM_PATH is defined and not set in ENV + elseif(DEFINED ROCM_PATH) + set(HIP_CLANG_PATH "${ROCM_PATH}/llvm/bin") else() - set(HIP_CLANG_PATH "/opt/rocm/llvm/bin") + message(FATAL_ERROR "Unable to find the clang compiler path. Set ROCM_PATH or HIP_PATH in env") endif() endif() set(CMAKE_HIP_LINK_EXECUTABLE "${HIP_HIPCC_CMAKE_LINKER_HELPER} ${HIP_CLANG_PATH} ${HIP_CLANG_PARALLEL_BUILD_LINK_OPTIONS} -o ") diff --git a/hip-lang-config.cmake.in b/hip-lang-config.cmake.in index dbc8e5348d..9250a68a07 100644 --- a/hip-lang-config.cmake.in +++ b/hip-lang-config.cmake.in @@ -65,8 +65,7 @@ find_dependency(amd_comgr) include( "${CMAKE_CURRENT_LIST_DIR}/hip-lang-targets.cmake" ) -#get_filename_component cannot resolve the symlinks if called from /opt/rocm/lib/hip -#and do three level up again +# From hip-lang config directory, do three level up get_filename_component(_DIR "${CMAKE_CURRENT_LIST_DIR}" REALPATH) get_filename_component(_IMPORT_PREFIX "${_DIR}/../../../" REALPATH) @@ -91,7 +90,6 @@ find_path(HSA_HEADER hsa/hsa.h "${_IMPORT_PREFIX}/../include" #FILE_REORG_BACKWARD_COMPATIBILITY "${_IMPORT_PREFIX}/include" "${ROCM_PATH}/include" - /opt/rocm/include ) if (NOT HSA_HEADER) From 08d24c55d9ec6fcf4d837af84b515446bc14c9d0 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Thu, 4 Aug 2022 09:32:12 +0530 Subject: [PATCH 6/9] Fix win:setenv/unsetenv tests build failure --- .../device/hipDeviceGetP2PAttribute_exe.cc | 34 +++++++++++++++---- 1 file changed, 27 insertions(+), 7 deletions(-) diff --git a/tests/catch/unit/device/hipDeviceGetP2PAttribute_exe.cc b/tests/catch/unit/device/hipDeviceGetP2PAttribute_exe.cc index 44e98495ca..92b1c577f5 100644 --- a/tests/catch/unit/device/hipDeviceGetP2PAttribute_exe.cc +++ b/tests/catch/unit/device/hipDeviceGetP2PAttribute_exe.cc @@ -20,23 +20,43 @@ THE SOFTWARE. #include #include "hip/hip_runtime_api.h" #include +#include +bool UNSETENV(std::string var) { + int result = -1; + #ifdef __unix__ + result = unsetenv(var.c_str()); + #else + result = _putenv((var + '=').c_str()); + #endif + return (result == 0) ? true: false; +} + +bool SETENV(std::string var, std::string value, int overwrite) { + int result = -1; + #ifdef __unix__ + result = setenv(var.c_str(), value.c_str(), overwrite); + #else + result = _putenv((var + '=' + value).c_str()); + #endif + return (result == 0) ? true: false; +} void inline hideDevices(const char* devices) { #if HT_NVIDIA - setenv("CUDA_VISIBLE_DEVICES", devices, 1); + SETENV("CUDA_VISIBLE_DEVICES", devices, 1); #else - setenv("HIP_VISIBLE_DEVICES", devices, 1); - setenv("ROCR_VISIBLE_DEVICES", devices, 1); + SETENV("HIP_VISIBLE_DEVICES", devices, 1); + SETENV("ROCR_VISIBLE_DEVICES", devices, 1); #endif } void inline unhideAllDevices() { #if HT_NVIDIA - unsetenv("CUDA_VISIBLE_DEVICES"); + UNSETENV("CUDA_VISIBLE_DEVICES"); #else - unsetenv("HIP_VISIBLE_DEVICES"); - unsetenv("ROCR_VISIBLE_DEVICES"); + UNSETENV("HIP_VISIBLE_DEVICES"); + UNSETENV("ROCR_VISIBLE_DEVICES"); #endif } @@ -63,4 +83,4 @@ int main(int argc, char** argv) { } return error; -} \ No newline at end of file +} From 04141114336a3fe707c98a1eab5d03a0f045c9f0 Mon Sep 17 00:00:00 2001 From: Anton Mitkov Date: Thu, 4 Aug 2022 06:05:21 +0100 Subject: [PATCH 7/9] EXSWCPHIPT-118 - Added testing for hipMemset Synchronous behavoiour. (#2750) --- tests/catch/include/hip_test_common.hh | 111 +++- tests/catch/unit/memory/CMakeLists.txt | 8 + tests/catch/unit/memory/DriverContext.hh | 2 + tests/catch/unit/memory/MemUtils.hh | 407 ++++++++++++++ tests/catch/unit/memory/hipFree.cc | 421 +++++++++++++++ tests/catch/unit/memory/hipMemGetInfo.cc | 233 ++++++-- tests/catch/unit/memory/hipMemcpySync.cc | 227 ++++++++ tests/catch/unit/memory/hipMemsetAsync.cc | 173 ++++++ tests/catch/unit/memory/hipMemsetSync.cc | 505 ++++++++++++++++++ .../unit/memory/hipPointerGetAttributes.cc | 164 +++--- .../hipStreamPerThread_Basic.cc | 4 +- 11 files changed, 2102 insertions(+), 153 deletions(-) create mode 100644 tests/catch/unit/memory/MemUtils.hh create mode 100644 tests/catch/unit/memory/hipFree.cc create mode 100644 tests/catch/unit/memory/hipMemcpySync.cc create mode 100644 tests/catch/unit/memory/hipMemsetAsync.cc create mode 100644 tests/catch/unit/memory/hipMemsetSync.cc diff --git a/tests/catch/include/hip_test_common.hh b/tests/catch/include/hip_test_common.hh index 0fb79a49b3..3a8501f62c 100644 --- a/tests/catch/include/hip_test_common.hh +++ b/tests/catch/include/hip_test_common.hh @@ -287,19 +287,19 @@ struct Pinned { //--- struct Unpinned { - static const bool isPinned = false; - static const char* str() { return "Unpinned"; }; + static const bool isPinned = false; + static const char* str() { return "Unpinned"; }; - static void* Alloc(size_t sizeBytes) { - void* p = malloc(sizeBytes); - HIPASSERT(p); - return p; - }; + static void* Alloc(size_t sizeBytes) { + void* p = malloc(sizeBytes); + HIPASSERT(p); + return p; + }; }; struct Memcpy { - static const char* str() { return "Memcpy"; }; + static const char* str() { return "Memcpy"; }; }; struct MemcpyAsync { @@ -307,33 +307,104 @@ struct MemcpyAsync { }; -template -struct MemTraits; +template struct MemTraits; -template <> -struct MemTraits { +template <> struct MemTraits { static void Copy(void* dest, const void* src, size_t sizeBytes, hipMemcpyKind kind, - hipStream_t stream) { + hipStream_t stream) { (void)stream; HIPCHECK(hipMemcpy(dest, src, sizeBytes, kind)); } }; -template <> -struct MemTraits { +template <> struct MemTraits { static void Copy(void* dest, const void* src, size_t sizeBytes, hipMemcpyKind kind, - hipStream_t stream) { + hipStream_t stream) { HIPCHECK(hipMemcpyAsync(dest, src, sizeBytes, kind, stream)); } }; -} // namespace HipTest +namespace { +static __global__ void waitKernel(clock_t offset) { + auto start = clock(); + while ((clock() - start) < offset) { + } +} + +// helper function used to set the device frequency variable +// estimates the number of clock ticks in 1 second +static size_t findTicksPerSecond() { + // first read the reported clockRate as a starting point + hipDeviceProp_t prop; + int device; + HIP_CHECK(hipGetDevice(&device)); + HIP_CHECK(hipGetDeviceProperties(&prop, device)); + clock_t devFreq = static_cast(prop.clockRate); // in kHz + clock_t clockTicksPerSecond = devFreq * 1000; + + // init + hipEvent_t start, stop; + HIP_CHECK(hipEventCreate(&start)); + HIP_CHECK(hipEventCreate(&stop)); + + // Warmup + hipLaunchKernelGGL(waitKernel, dim3(1), dim3(1), 0, 0, clockTicksPerSecond); + HIP_CHECK(hipGetLastError()); + HIP_CHECK(hipDeviceSynchronize()); + + // try 10 times to find device frequency + // after 10 attempts the result is likely good enough so just accept it + for (int attempts = 10; attempts > 0; --attempts) { + HIP_CHECK(hipEventRecord(start)); + hipLaunchKernelGGL(waitKernel, dim3(1), dim3(1), 0, 0, clockTicksPerSecond); + HIP_CHECK(hipEventRecord(stop)); + HIP_CHECK(hipGetLastError()); + HIP_CHECK(hipEventSynchronize(stop)); + + float executionTimeMs = 0; + HIP_CHECK(hipEventElapsedTime(&executionTimeMs, start, stop)); + + constexpr float tolerance = 20; + if (fabs(executionTimeMs - 1000) <= tolerance) { + // Timing is within accepted tolerance, break here + break; + } else { + clockTicksPerSecond = (clockTicksPerSecond * 1000) / executionTimeMs; + --attempts; + } + } + + // deinit + HIP_CHECK(hipEventDestroy(start)); + HIP_CHECK(hipEventDestroy(stop)); + return clockTicksPerSecond; +} +} // namespace + +// Launches a kernel which runs for specified amount of time +// Note: The current implementation uses HIP_CHECK which is not thread safe! +// Note: the function assumes execution on a single device and caches the number of clock ticks per +// second +static inline void runKernelForDuration(std::chrono::milliseconds duration, + hipStream_t stream = nullptr) { + // number of clocks the device is running at (device frequency) + // each translation unit will have a copy of ticksPerSecond but this function isn't designed for + // precision so that's acceptable. + static size_t ticksPerSecond = findTicksPerSecond(); + const auto millis = duration.count(); + hipLaunchKernelGGL(waitKernel, dim3(1), dim3(1), 0, stream, ticksPerSecond * millis / 1000); + HIP_CHECK(hipGetLastError()); +} + +} // namespace HipTest // This must be called in the beginning of image test app's main() to indicate whether image // is supported. -#define CHECK_IMAGE_SUPPORT \ - if (!HipTest::isImageSupported()) \ - { INFO("Texture is not support on the device. Skipped."); return; } +#define CHECK_IMAGE_SUPPORT \ + if (!HipTest::isImageSupported()) { \ + INFO("Texture is not support on the device. Skipped."); \ + return; \ + } diff --git a/tests/catch/unit/memory/CMakeLists.txt b/tests/catch/unit/memory/CMakeLists.txt index 853c1453af..ce52b2e104 100644 --- a/tests/catch/unit/memory/CMakeLists.txt +++ b/tests/catch/unit/memory/CMakeLists.txt @@ -92,6 +92,10 @@ set(TEST_SRC hipArray.cc hipMemVmm.cc hipMemGetInfo.cc + hipFree.cc + hipMemcpySync.cc + hipMemsetSync.cc + hipMemsetAsync.cc ) else() set(TEST_SRC @@ -159,6 +163,10 @@ set(TEST_SRC hipDrvPtrGetAttributes.cc hipMemPrefetchAsync.cc hipMemGetInfo.cc + hipFree.cc + hipMemcpySync.cc + hipMemsetSync.cc + hipMemsetAsync.cc ) endif() diff --git a/tests/catch/unit/memory/DriverContext.hh b/tests/catch/unit/memory/DriverContext.hh index 5593c512d4..24a84dcd3d 100644 --- a/tests/catch/unit/memory/DriverContext.hh +++ b/tests/catch/unit/memory/DriverContext.hh @@ -21,6 +21,8 @@ THE SOFTWARE. */ #pragma once +#include +#include #include diff --git a/tests/catch/unit/memory/MemUtils.hh b/tests/catch/unit/memory/MemUtils.hh new file mode 100644 index 0000000000..a05588c3b7 --- /dev/null +++ b/tests/catch/unit/memory/MemUtils.hh @@ -0,0 +1,407 @@ +/* +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 + + +namespace mem_utils { + +enum class allocType { deviceMalloc, hostMalloc, hostRegisted, devRegistered }; +enum class memType { hipMem, hipMemsetD8, hipMemsetD16, hipMemsetD32, hipMem2D, hipMem3D }; + +// helper struct containing vars needed for 2D and 3D mem Testing +struct MultiDData { + size_t width{}; // in elements not bytes + // set to 0 for 1D + size_t height{}; // in elements not bytes + size_t getH() { return height == 0 ? 1 : height; }; // return 1 if height == 0 || height + // set to 0 for 2D + size_t depth{}; // in elements not bytes + size_t getD() { return depth == 0 ? 1 : depth; }; // return 1 if depth == 0 || depth + size_t pitch{}; // pitch = (width * sizeofData) + alignment + size_t offset{}; // for simplicity use same offset for x,y and z dimentions of memory + size_t getCount() { return width * getH() * getD(); } +}; + +// set of helper functions to tidy the nested switch statements +template +static inline std::pair deviceMallocHelper(memType memType, size_t dataW, size_t dataH, + size_t dataD, size_t& dataPitch) { + constexpr size_t elementSize = sizeof(T); + size_t sizeInBytes = elementSize * dataW * dataH * dataD; + T* aPtr{}; + switch (memType) { + case memType::hipMem3D: { + hipPitchedPtr pitchedAPtr; + hipExtent extent = make_hipExtent(dataW * elementSize, dataH, dataD); + + HIP_CHECK(hipMalloc3D(&pitchedAPtr, extent)); + aPtr = reinterpret_cast(pitchedAPtr.ptr); + dataPitch = pitchedAPtr.pitch; + break; + } + + case memType::hipMem2D: + HIP_CHECK( + hipMallocPitch(reinterpret_cast(&aPtr), &dataPitch, dataW * elementSize, dataH)); + + break; + + default: + HIP_CHECK(hipMalloc(&aPtr, sizeInBytes)); + dataPitch = dataW * elementSize; + break; + } + return {aPtr, nullptr}; +} + +template +static inline std::pair hostMallocHelper(size_t dataW, size_t dataH, size_t dataD, + size_t& dataPitch) { + constexpr size_t elementSize = sizeof(T); + size_t sizeInBytes = elementSize * dataW * dataH * dataD; + T* aPtr; + + HIP_CHECK(hipHostMalloc(&aPtr, sizeInBytes)); + dataPitch = dataW * elementSize; + + return {aPtr, nullptr}; +} + +template +static inline std::pair hostRegisteredHelper(size_t dataW, size_t dataH, size_t dataD, + size_t& dataPitch) { + constexpr size_t elementSize = sizeof(T); + size_t sizeInBytes = elementSize * dataW * dataH * dataD; + T* aPtr = new T[dataW * dataH * dataD]; + + HIP_CHECK(hipHostRegister(aPtr, sizeInBytes, hipHostRegisterDefault)); + + dataPitch = dataW * elementSize; + return {aPtr, nullptr}; +} + +template +static inline std::pair devRegisteredHelper(size_t dataW, size_t dataH, size_t dataD, + size_t& dataPitch) { + constexpr size_t elementSize = sizeof(T); + size_t sizeInBytes = elementSize * dataW * dataH * dataD; + T* aPtr = new T[dataW * dataH * dataD]; + T* retPtr{}; + + HIP_CHECK(hipHostRegister(aPtr, sizeInBytes, hipHostRegisterDefault)); + HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&retPtr), aPtr, 0)); + + dataPitch = dataW * elementSize; + // keep the address of the host memory + return {retPtr, aPtr}; +} + +/* + * helper function to allocate memory and set it to a value. + * return a pair of pointers due to the device registered allocation case, we need to keep track of + * the pointer to host memory to be able to unregister and free it + */ +template +static inline std::pair initMemory(allocType type, memType memType, MultiDData& data) { + std::pair retPtr{}; + // check different types of allocation + switch (type) { + case allocType::deviceMalloc: + retPtr = deviceMallocHelper(memType, data.width, data.getH(), data.getD(), data.pitch); + break; + + case allocType::hostMalloc: + retPtr = hostMallocHelper(data.width, data.getH(), data.getD(), data.pitch); + break; + + case allocType::hostRegisted: + retPtr = hostRegisteredHelper(data.width, data.getH(), data.getD(), data.pitch); + break; + + case allocType::devRegistered: + retPtr = devRegisteredHelper(data.width, data.getH(), data.getD(), data.pitch); + break; + + default: + REQUIRE(false); + break; + } + return retPtr; +} +// create a hipMemcpy3DParams struct for the 3d version of memcpy to verify the memset operation +template +hipMemcpy3DParms createParams(hipMemcpyKind kind, T* src, T* host_dst, size_t srcPitch, + size_t dataW, size_t dataH, size_t dataD) { + hipMemcpy3DParms p = {}; + p.kind = kind; + + p.srcPtr.ptr = src; + p.srcPtr.pitch = srcPitch; + p.srcPtr.xsize = dataW; + p.srcPtr.ysize = dataH; + + p.dstPtr.ptr = host_dst; + p.dstPtr.pitch = dataW * sizeof(T); + p.dstPtr.xsize = dataW; + p.dstPtr.ysize = dataH; + + hipExtent extent = make_hipExtent(dataW * sizeof(T), dataH, dataD); + p.extent = extent; + + return p; +} + +// set of helper functions to tidy the nested switch statements +template +static inline void deviceMallocCopy(memType memType, T* aPtr, T* hostMem, size_t dataW, + size_t dataH, size_t dataD, size_t& dataPitch) { + constexpr size_t elementSize = sizeof(T); + size_t sizeInBytes = elementSize * dataW * dataH * dataD; + switch (memType) { + case memType::hipMem3D: { + hipMemcpy3DParms params = + createParams(hipMemcpyDeviceToHost, aPtr, hostMem, dataPitch, dataW, dataH, dataD); + HIP_CHECK(hipMemcpy3D(¶ms)); + break; + } + + case memType::hipMem2D: + HIP_CHECK(hipMemcpy2D(hostMem, dataW * elementSize, aPtr, dataPitch, dataW, dataH, + hipMemcpyDeviceToHost)); + break; + + default: + HIP_CHECK(hipMemcpy(hostMem, aPtr, sizeInBytes, hipMemcpyDeviceToHost)); + break; + } +} + +template +static inline void hostCopy(memType memType, T* aPtr, T* hostMem, size_t dataW, size_t dataH, + size_t dataD, size_t& dataPitch) { + constexpr size_t elementSize = sizeof(T); + size_t sizeInBytes = elementSize * dataW * dataH * dataD; + switch (memType) { + case memType::hipMem3D: { + hipMemcpy3DParms params = + createParams(hipMemcpyHostToHost, aPtr, hostMem, dataPitch, dataW, dataH, dataD); + + HIP_CHECK(hipMemcpy3D(¶ms)); + break; + } + + case memType::hipMem2D: + HIP_CHECK(hipMemcpy2D(hostMem, dataW * elementSize, aPtr, dataPitch, dataW, dataH, + hipMemcpyHostToHost)); + break; + + default: + HIP_CHECK(hipMemcpy(hostMem, aPtr, sizeInBytes, hipMemcpyHostToHost)); + break; + } +} + +template +static inline void devRegisteredCopy(memType memType, T* aPtr, T* hostMem, size_t dataW, + size_t dataH, size_t dataD, size_t& dataPitch) { + constexpr size_t elementSize = sizeof(T); + + switch (memType) { + case memType::hipMem3D: { + hipMemcpy3DParms params = + createParams(hipMemcpyDeviceToHost, aPtr, hostMem, dataPitch, dataW, dataH, dataD); + + HIP_CHECK(hipMemcpy3D(¶ms)); + break; + } + + case memType::hipMem2D: + HIP_CHECK(hipMemcpy2D(hostMem, dataW * elementSize, aPtr, dataPitch, dataW, dataH, + hipMemcpyDeviceToHost)); + break; + + default: { + size_t sizeInBytes = elementSize * dataW * dataH * dataD; + HIP_CHECK(hipMemcpy(hostMem, aPtr, sizeInBytes, hipMemcpyDeviceToHost)); + break; + } + } +} + +/* + * function returns an offset location in memory based on the provided data, taking pitch into + * account + * (for 1D requires data.depth & data.height = 0, for 2D data.depth = 0) + */ +static inline size_t getPtrOffset(MultiDData data) { + if (data.height == 0) { // 1D + return data.offset; + } else if (data.depth == 0) { + return (data.offset + (data.pitch * data.offset)); + } else { // 2D or 3D + return (data.offset + (data.pitch * data.offset) + (data.pitch * data.offset * data.height)); + } +} + +/* + * Function to allow reuse of functions for testing versions of the memset API, at a specified + * offset + */ +template +static inline void memsetCheck(T* aPtr, size_t value, memType memType, MultiDData& data, + hipStream_t stream = nullptr, bool async = true) { + size_t count = data.getCount(); + size_t ptrOffset{}; + switch (memType) { + case memType::hipMem: + if (async) { + HIP_CHECK(hipMemsetAsync(aPtr + data.offset, value, count * sizeof(T), stream)); + } else { + HIP_CHECK(hipMemset(aPtr + data.offset, value, count * sizeof(T))); + } + break; + + case memType::hipMemsetD8: + if (async) { + HIP_CHECK(hipMemsetD8Async(reinterpret_cast(aPtr + data.offset), value, + count, stream)); + } else { + HIP_CHECK(hipMemsetD8(reinterpret_cast(aPtr + data.offset), value, count)); + } + break; + + case memType::hipMemsetD16: + if (async) { + HIP_CHECK(hipMemsetD16Async(reinterpret_cast(aPtr + data.offset), value, + count, stream)); + } else { + HIP_CHECK(hipMemsetD16(reinterpret_cast(aPtr + data.offset), value, count)); + } + break; + + case memType::hipMemsetD32: + if (async) { + HIP_CHECK(hipMemsetD32Async(reinterpret_cast(aPtr + data.offset), value, + count, stream)); + } else { + HIP_CHECK(hipMemsetD32(reinterpret_cast(aPtr + data.offset), value, count)); + } + break; + + case memType::hipMem2D: + ptrOffset = getPtrOffset(data); + if (async) { + HIP_CHECK( + hipMemset2DAsync(aPtr + ptrOffset, data.pitch, value, data.width, data.height, stream)); + } else { + HIP_CHECK(hipMemset2D(aPtr + ptrOffset, data.pitch, value, data.width, data.height)); + } + break; + + case memType::hipMem3D: { + ptrOffset = getPtrOffset(data); + hipExtent extent = make_hipExtent(data.width * sizeof(T), data.height, data.depth); + + if (async) { + HIP_CHECK(hipMemset3DAsync( + make_hipPitchedPtr(aPtr + ptrOffset, data.pitch, data.width, data.height), value, + extent, stream)); + } else { + HIP_CHECK( + hipMemset3D(make_hipPitchedPtr(aPtr + ptrOffset, data.pitch, data.width, data.height), + value, extent)); + } + break; + } + default: + REQUIRE(false); + break; + } +} + +template static inline void freeStuff(T* aPtr, allocType type) { + switch (type) { + case allocType::deviceMalloc: + hipFree(aPtr); + break; + case allocType::hostMalloc: + hipHostFree(aPtr); + break; + default: // for host and device registered + HIP_CHECK(hipHostUnregister(aPtr)); + delete[] aPtr; + break; + } +} + +/* + * Copies device data to host and checks that each element is equal to the + * specified value + */ +template +static inline void verifyData(T* aPtr, size_t value, MultiDData& data, allocType type, + memType memType) { + std::unique_ptr hostPtr = std::make_unique(data.getCount()); + switch (type) { + case allocType::deviceMalloc: + deviceMallocCopy(memType, aPtr + getPtrOffset(data), hostPtr.get(), data.width, data.getH(), + data.getD(), data.pitch); + break; + case allocType::devRegistered: + devRegisteredCopy(memType, aPtr + getPtrOffset(data), hostPtr.get(), data.width, data.getH(), + data.getD(), data.pitch); + break; + default: // host malloc and host registered + hostCopy(memType, aPtr + getPtrOffset(data), hostPtr.get(), data.width, data.getH(), + data.getD(), data.pitch); + break; + } + + size_t idx; + bool allMatch{true}; + for (size_t k = 0; k < data.getD(); k++) { + for (size_t j = 0; j < data.getH(); j++) { + for (size_t i = 0; i < data.width; i++) { + idx = data.width * data.getH() * k + data.width * j + i; + allMatch = allMatch && static_cast(hostPtr.get()[idx]) == value; + if (!allMatch) REQUIRE(false); + } + } + } +} + +// function used to abstract the test +template +static inline void doMemTest(F func, fArgs... funcArgs) { + SECTION("Synchronous") { func(nullptr, false, funcArgs...); } + SECTION("Asynchronous - null stream") { func(nullptr, true, funcArgs...); } + SECTION("Asynchronous - created stream") { + hipStream_t stream{}; + HIP_CHECK(hipStreamCreate(&stream)); + func(stream, true, funcArgs...); + HIP_CHECK(hipStreamDestroy(stream)); + } +} +} // namespace mem_utils diff --git a/tests/catch/unit/memory/hipFree.cc b/tests/catch/unit/memory/hipFree.cc new file mode 100644 index 0000000000..1248deebc1 --- /dev/null +++ b/tests/catch/unit/memory/hipFree.cc @@ -0,0 +1,421 @@ +/* +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 "hipArrayCommon.hh" +#include "DriverContext.hh" + +/* + * This testcase verifies [ hipFree || hipFreeArray || hipFreeType::ArrayDestroy || + * hipFreeType::HostFree with hipHostMalloc ] + * 1. Check that hipFree implicitly synchronises the device. + * 2. Perform multiple allocations and then call hipFree on each pointer concurrently (from unique + * threads) for different memory types and different allocation sizes. + * 3. Pass nullptr as argument and check that no operation is performed and hipSuccess is returned. + * 4. Pass an invalid ptr and check that hipErrorInvalidValue is returned. + * 5. Call hipFree twice on the same pointer and check that the implementation handles the second + * call correctly. + * 6. HipFreeType::HostFree only: + * Try to free memory that has been registered with hipHostRegister and check that + * hipErrorInvalidValue is returned. + */ + + +enum class FreeType { DevFree, ArrayFree, ArrayDestroy, HostFree }; + +// Amount of time kernel should wait +using namespace std::chrono_literals; +const std::chrono::duration delay = 50ms; +constexpr size_t numAllocs = 10; + +#if HT_AMD /* Disabled because frequency based wait is timing out on nvidia platforms */ +TEMPLATE_TEST_CASE("Unit_hipFreeImplicitSyncDev", "", char, float, float2, float4) { + TestType* devPtr{}; + size_t size_mult = GENERATE(1, 32, 64, 128, 256); + HIP_CHECK(hipMalloc(&devPtr, sizeof(TestType) * size_mult)); + + HipTest::runKernelForDuration(delay); + // make sure device is busy + HIP_CHECK_ERROR(hipStreamQuery(nullptr), hipErrorNotReady); + HIP_CHECK(hipFree(devPtr)); + HIP_CHECK(hipStreamQuery(nullptr)); +} + +TEMPLATE_TEST_CASE("Unit_hipFreeImplicitSyncHost", "", char, float, float2, float4) { + TestType* hostPtr{}; + size_t size_mult = GENERATE(1, 32, 64, 128, 256); + + HIP_CHECK(hipHostMalloc(&hostPtr, sizeof(TestType) * size_mult)); + + HipTest::runKernelForDuration(delay); + // make sure device is busy + HIP_CHECK_ERROR(hipStreamQuery(nullptr), hipErrorNotReady); + HIP_CHECK(hipHostFree(hostPtr)); + HIP_CHECK(hipStreamQuery(nullptr)); +} + +#if HT_NVIDIA // Meaningless at the moment, since we are not running wait kernel on nvidia. +TEMPLATE_TEST_CASE("Unit_hipFreeImplicitSyncArray", "", char, float, float2, float4) { + using vec_info = vector_info; + DriverContext ctx; + + + size_t width = GENERATE(32, 512, 1024); + size_t height = GENERATE(32, 512, 1024); + + SECTION("ArrayFree") { + hipArray_t arrayPtr{}; + hipChannelFormatDesc desc = hipCreateChannelDesc(); + + HIP_CHECK(hipMallocArray(&arrayPtr, &desc, width, height, hipArrayDefault)); + HipTest::runKernelForDuration(delay); + // make sure device is busy + HIP_CHECK_ERROR(hipStreamQuery(nullptr), hipErrorNotReady); + HIP_CHECK(hipFreeArray(arrayPtr)); + HIP_CHECK(hipStreamQuery(nullptr)); + } + SECTION("ArrayDestroy") { + hiparray cuArrayPtr{}; + + HIP_ARRAY_DESCRIPTOR cuDesc; + cuDesc.Width = width; + cuDesc.Height = height; + cuDesc.Format = vec_info::format; + cuDesc.NumChannels = vec_info::size; + HIP_CHECK(hipArrayCreate(&cuArrayPtr, &cuDesc)); + HipTest::runKernelForDuration(delay); + // make sure device is busy + HIP_CHECK_ERROR(hipStreamQuery(nullptr), hipErrorNotReady); + HIP_CHECK(hipArrayDestroy(cuArrayPtr)); + HIP_CHECK(hipStreamQuery(nullptr)); + } +} +#else // AMD + +TEMPLATE_TEST_CASE("Unit_hipFreeImplicitSyncArray", "", char, float, float2, float4) { + hipArray_t arrayPtr{}; + hipExtent extent{}; + extent.width = GENERATE(32, 128, 256, 512, 1024); + extent.height = GENERATE(0, 32, 128, 256, 512, 1024); + hipChannelFormatDesc desc = hipCreateChannelDesc(); + + HIP_CHECK(hipMallocArray(&arrayPtr, &desc, extent.width, extent.height, hipArrayDefault)); + HipTest::runKernelForDuration(delay); + // make sure device is busy + HIP_CHECK_ERROR(hipStreamQuery(nullptr), hipErrorNotReady); + // Second free segfaults + SECTION("ArrayDestroy") { + HIP_CHECK(hipArrayDestroy(arrayPtr)); + HIP_CHECK(hipStreamQuery(nullptr)); + } + SECTION("ArrayFree") { + HIP_CHECK(hipFreeArray(arrayPtr)); + HIP_CHECK(hipStreamQuery(nullptr)); + } +} + +#endif +#endif + +// Freeing a invalid pointer with on device +TEST_CASE("Unit_hipFreeNegativeDev") { + SECTION("InvalidPtr") { + char value; + HIP_CHECK_ERROR(hipFree(&value), hipErrorInvalidValue); + } + SECTION("NullPtr") { HIP_CHECK(hipFree(nullptr)); } +} + +// Freeing a invalid pointer with on host +TEST_CASE("Unit_hipFreeNegativeHost") { + SECTION("NullPtr") { HIP_CHECK(hipHostFree(nullptr)); } + SECTION("InvalidPtr") { + char hostPtr; + HIP_CHECK_ERROR(hipHostFree(&hostPtr), hipErrorInvalidValue); + } + SECTION("hipHostRegister") { + char* hostPtr = new char; + auto flag = GENERATE(hipHostRegisterDefault, hipHostRegisterPortable, hipHostRegisterMapped); + HIP_CHECK(hipHostRegister((void*)hostPtr, sizeof(char), flag)); + HIP_CHECK_ERROR(hipHostFree(hostPtr), hipErrorInvalidValue); + delete hostPtr; + } +} + +#if HT_NVIDIA +TEST_CASE("Unit_hipFreeNegativeArray") { + DriverContext ctx; + hipArray_t arrayPtr{}; + hiparray cuArrayPtr{}; + + SECTION("ArrayFree") { HIP_CHECK(hipFreeArray(nullptr)); } + SECTION("ArrayDestroy") { + HIP_CHECK_ERROR(hipArrayDestroy(nullptr), hipErrorInvalidResourceHandle); + } +} +#else + +// Freeing a invalid pointer with array +TEST_CASE("Unit_hipFreeNegativeArray") { + SECTION("ArrayFree") { HIP_CHECK_ERROR(hipFreeArray(nullptr), hipErrorInvalidValue); } + SECTION("ArrayDestroy") { HIP_CHECK_ERROR(hipArrayDestroy(nullptr), hipErrorInvalidValue); } +} + +#endif + +TEST_CASE("Unit_hipFreeDoubleDevice") { + size_t width = GENERATE(32, 512, 1024); + char* ptr{}; + size_t size_mult = width; + HIP_CHECK(hipMalloc(&ptr, sizeof(char) * size_mult)); + + HIP_CHECK(hipFree(ptr)); + HIP_CHECK_ERROR(hipFree(ptr), hipErrorInvalidValue); +} +TEST_CASE("Unit_hipFreeDoubleHost") { + size_t width = GENERATE(32, 512, 1024); + char* ptr{}; + size_t size_mult = width; + + HIP_CHECK(hipHostMalloc(&ptr, sizeof(char) * size_mult)); + + HIP_CHECK(hipHostFree(ptr)); + HIP_CHECK_ERROR(hipHostFree(ptr), hipErrorInvalidValue); +} + +#if HT_NVIDIA +TEST_CASE("Unit_hipFreeDoubleArrayFree") { + HipTest::HIP_SKIP_TEST("EXSWCPHIPT-120"); + return; + + size_t width = GENERATE(32, 512, 1024); + size_t height = GENERATE(0, 32, 512, 1024); + hipArray_t arrayPtr{}; + hipExtent extent{}; + extent.width = width; + extent.height = height; + hipChannelFormatDesc desc = hipCreateChannelDesc(); + + HIP_CHECK(hipMallocArray(&arrayPtr, &desc, extent.width, extent.height, hipArrayDefault)); + + HIP_CHECK(hipFreeArray(arrayPtr)); + HIP_CHECK_ERROR(hipFreeArray(arrayPtr), hipErrorContextIsDestroyed); +} + +TEST_CASE("Unit_hipFreeDoubleArrayDestroy") { + HipTest::HIP_SKIP_TEST("EXSWCPHIPT-120"); + return; + using vec_info = vector_info; + + size_t width = GENERATE(32, 512, 1024); + size_t height = GENERATE(0, 32, 512, 1024); + DriverContext ctx{}; + + hiparray ArrayPtr{}; + HIP_ARRAY_DESCRIPTOR cuDesc; + cuDesc.Width = width; + cuDesc.Height = height; + cuDesc.Format = vec_info::format; + cuDesc.NumChannels = vec_info::size; + HIP_CHECK(hipArrayCreate(&ArrayPtr, &cuDesc)); + HIP_CHECK(hipArrayDestroy(ArrayPtr)); + HIP_CHECK_ERROR(hipArrayDestroy(ArrayPtr), hipErrorContextIsDestroyed); +} + +#else // AMD + +TEST_CASE("Unit_hipFreeDoubleArray") { + size_t width = GENERATE(32, 512, 1024); + size_t height = GENERATE(0, 32, 512, 1024); + hipArray_t arrayPtr{}; + hipExtent extent{}; + extent.width = width; + extent.height = height; + hipChannelFormatDesc desc = hipCreateChannelDesc(); + + HIP_CHECK(hipMallocArray(&arrayPtr, &desc, extent.width, extent.height, hipArrayDefault)); + + SECTION("ArrayFree") { + HIP_CHECK(hipFreeArray(arrayPtr)); + HIP_CHECK_ERROR(hipFreeArray(arrayPtr), hipErrorContextIsDestroyed); + } + SECTION("ArrayDestroy") { + HIP_CHECK(hipArrayDestroy(arrayPtr)); + HIP_CHECK_ERROR(hipArrayDestroy(arrayPtr), hipErrorContextIsDestroyed); + } +} + +#endif + + +TEMPLATE_TEST_CASE("Unit_hipFreeMultiTDev", "", char, int, float2, float4) { + std::vector ptrs(numAllocs); + size_t allocSize = sizeof(TestType) * GENERATE(1, 32, 64, 128); + + for (auto& ptr : ptrs) { + HIP_CHECK(hipMalloc(&ptr, allocSize)); + } + + std::vector threads; + + for (auto ptr : ptrs) { + threads.emplace_back(([ptr] { + HIP_CHECK_THREAD(hipFree(ptr)); + HIP_CHECK_THREAD(hipStreamQuery(nullptr)); + })); + } + + for (auto& t : threads) { + t.join(); + } + HIP_CHECK_THREAD_FINALIZE(); +} + +TEMPLATE_TEST_CASE("Unit_hipFreeMultiTHost", "", char, int, float2, float4) { + std::vector ptrs(numAllocs); + size_t allocSize = sizeof(TestType) * GENERATE(1, 32, 64, 128); + + for (auto& ptr : ptrs) { + HIP_CHECK(hipHostMalloc(&ptr, allocSize)); + } + + std::vector threads; + + for (auto ptr : ptrs) { + threads.emplace_back(([ptr] { + HIP_CHECK_THREAD(hipHostFree(ptr)); + HIP_CHECK_THREAD(hipStreamQuery(nullptr)); + })); + } + + for (auto& t : threads) { + t.join(); + } + HIP_CHECK_THREAD_FINALIZE(); +} + +#if HT_NVIDIA +TEMPLATE_TEST_CASE("Unit_hipFreeMultiTArray", "", char, int, float2, float4) { + using vec_info = vector_info; + + size_t width = GENERATE(32, 128, 256, 512, 1024); + size_t height = GENERATE(32, 128, 256, 512, 1024); + DriverContext ctx; + std::vector threads; + + + SECTION("ArrayDestroy") { + std::vector ptrs(numAllocs); + HIP_ARRAY_DESCRIPTOR cuDesc; + cuDesc.Width = width; + cuDesc.Height = height; + cuDesc.Format = vec_info::format; + cuDesc.NumChannels = vec_info::size; + for (auto& ptr : ptrs) { + HIP_CHECK(hipArrayCreate(&ptr, &cuDesc)); + } + + + for (auto& ptr : ptrs) { + threads.emplace_back(([ptr] { + HIP_CHECK_THREAD(hipArrayDestroy(ptr)); + HIP_CHECK_THREAD(hipStreamQuery(nullptr)); + })); + } + for (auto& t : threads) { + t.join(); + } + HIP_CHECK_THREAD_FINALIZE(); + } + + SECTION("ArrayFree") { + std::vector ptrs(numAllocs); + hipExtent extent{}; + extent.width = width; + extent.height = height; + hipChannelFormatDesc desc = hipCreateChannelDesc(); + + for (auto& ptr : ptrs) { + HIP_CHECK(hipMallocArray(&ptr, &desc, extent.width, extent.height, hipArrayDefault)); + } + + for (auto ptr : ptrs) { + SECTION("ArrayFree") { + threads.emplace_back(([ptr] { + HIP_CHECK_THREAD(hipFreeArray(ptr)); + HIP_CHECK_THREAD(hipStreamQuery(nullptr)); + })); + } + } + for (auto& t : threads) { + t.join(); + } + HIP_CHECK_THREAD_FINALIZE(); + } +} +#else + +TEMPLATE_TEST_CASE("Unit_hipFreeMultiTArray", "", char, int, float2, float4) { + using vec_info = vector_info; + + hipExtent extent{}; + extent.width = GENERATE(32, 128, 256, 512, 1024); + extent.height = GENERATE(0, 32, 128, 256, 512, 1024); + hipChannelFormatDesc desc = hipCreateChannelDesc(); + + std::vector threads; + + SECTION("ArrayFree") { + std::vector ptrs(numAllocs); + for (auto& ptr : ptrs) { + HIP_CHECK(hipMallocArray(&ptr, &desc, extent.width, extent.height, hipArrayDefault)); + threads.emplace_back([ptr] { + HIP_CHECK_THREAD(hipFreeArray(ptr)); + HIP_CHECK_THREAD(hipStreamQuery(nullptr)); + }); + } + } + SECTION("ArrayDestroy") { + std::vector cuArrayPtrs(numAllocs); + + HIP_ARRAY_DESCRIPTOR cuDesc; + cuDesc.Width = extent.width; + cuDesc.Height = extent.height; + cuDesc.Format = vec_info::format; + cuDesc.NumChannels = vec_info::size; + for (auto ptr : cuArrayPtrs) { + HIP_CHECK(hipArrayCreate(&ptr, &cuDesc)); + + threads.emplace_back([ptr] { + HIP_CHECK_THREAD(hipArrayDestroy(ptr)); + HIP_CHECK_THREAD(hipStreamQuery(nullptr)); + }); + } + } + for (auto& t : threads) { + t.join(); + } + HIP_CHECK_THREAD_FINALIZE(); +} + +#endif \ No newline at end of file diff --git a/tests/catch/unit/memory/hipMemGetInfo.cc b/tests/catch/unit/memory/hipMemGetInfo.cc index fc91511a3b..93481eb347 100644 --- a/tests/catch/unit/memory/hipMemGetInfo.cc +++ b/tests/catch/unit/memory/hipMemGetInfo.cc @@ -1,16 +1,13 @@ /* 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 @@ -28,12 +25,13 @@ THE SOFTWARE. /* * This testcase verifies hipMemGetInfo API * 1. Different memory chunk allocation - * 1.1. hipMalloc + * 1.1. hipMalloc - smallest memory chunck that can be allocated is 1024 * 1.2. hipMallocArray * 1.3. hipMalloc3D * 1.3. hipMalloc3DArray * 2. Allocation using different threads * 3. Negative: Invalid args + * */ struct MinAlloc { @@ -71,9 +69,9 @@ struct MinAlloc { // if the memory being allocated is not divisible by the minimum allocation add an extra minimum // allocation AddedAllocation = InitialAllocation + (MinAllocation - divisionRemainer) void fixAllocSize(size_t& allocation) { - REQUIRE(MinAlloc::Get() != 0); + REQUIRE(MinAlloc::Get() >= 0); if (allocation % MinAlloc::Get() != 0) { - auto adjustment = allocation % MinAlloc::Get(); + auto adjustment = allocation % MinAlloc::Get(); // FIXME This does mod by zero adjustment = MinAlloc::Get() - adjustment; allocation = allocation + adjustment; } @@ -88,6 +86,48 @@ void fixAllocSize(size_t& allocation) { << "Memory assumed to be used: \t\t" << usedMem); +TEST_CASE("Unit_hipMemGetInfo_DifferentMallocSmall") { + size_t freeMemInit; + size_t totalMemInit; + HIP_CHECK(hipMemGetInfo(&freeMemInit, &totalMemInit)); + + unsigned int* A_mem{nullptr}; + size_t freeMemRet; + size_t totalMemRet; + // allocate smaller chunk than minimum + size_t Malloc1Size = 2; + + HIP_CHECK(hipMalloc(reinterpret_cast(&A_mem), Malloc1Size)); + + HIP_CHECK(hipMemGetInfo(&freeMemRet, &totalMemRet)); + MEMINFO(totalMemRet, freeMemInit, freeMemRet, Malloc1Size); + + auto assumedFreeMem = freeMemInit - Malloc1Size; + // Free memory should be less than assumed for + // single allocation smaller than min allocation chunk + REQUIRE(freeMemRet < assumedFreeMem); + // confirms that allocated memory is at least equal to smallest allocation + assumedFreeMem = freeMemInit - MinAlloc::Get(); + REQUIRE(freeMemRet <= assumedFreeMem); + + HIP_CHECK(hipFree(A_mem)); + + // allocate smallest chunk of memory + HIP_CHECK(hipMalloc(reinterpret_cast(&A_mem), MinAlloc::Get())); + HIP_CHECK(hipMemGetInfo(&freeMemRet, &totalMemRet)); + + MEMINFO(totalMemRet, freeMemInit, freeMemRet, MinAlloc::Get()); + + assumedFreeMem = freeMemInit - MinAlloc::Get(); + // confirms that allocated memory is at least equal to smallest allocation + REQUIRE(freeMemRet <= assumedFreeMem); + + HIP_CHECK(hipFree(A_mem)); +} + +#if 0 // FIXME_jatinx Disabled for now because the formula to calulcate memget info is incorrect + // To be enabled after correct formula is found. + TEST_CASE("Unit_hipMemGetInfo_DifferentMallocLarge") { size_t freeMemInit; size_t totalMemInit; @@ -109,10 +149,12 @@ TEST_CASE("Unit_hipMemGetInfo_DifferentMallocLarge") { auto Malloc1Size = freeMemInit >> 1; // if the allocation is not divisible by the MinAllocation // take into account and add padding + fixAllocSize(Malloc1Size); HIP_CHECK(hipMalloc(reinterpret_cast(&A_mem), Malloc1Size)); // allocate an extra quarter of free mem auto Malloc2Size = Malloc1Size >> 1; + fixAllocSize(Malloc2Size); HIP_CHECK(hipMalloc(reinterpret_cast(&B_mem), Malloc2Size)); HIP_CHECK(hipMemGetInfo(&freeMemRet, &totalMemRet)); @@ -129,29 +171,6 @@ TEST_CASE("Unit_hipMemGetInfo_DifferentMallocLarge") { HIP_CHECK(hipFree(B_mem)); } -TEST_CASE("Unit_hipMemGetInfo_DifferentMallocSmall") { - size_t freeMemInit; - size_t totalMemInit; - HIP_CHECK(hipMemGetInfo(&freeMemInit, &totalMemInit)); - - unsigned int* A_mem{nullptr}; - size_t freeMemRet; - size_t totalMemRet; - // allocate smaller chunk than minimum - size_t Malloc1Size = 1; - - HIP_CHECK(hipMalloc(reinterpret_cast(&A_mem), Malloc1Size)); - - HIP_CHECK(hipMemGetInfo(&freeMemRet, &totalMemRet)); - MEMINFO(totalMemRet, freeMemInit, freeMemRet, Malloc1Size); - - auto assumedFreeMem = freeMemInit - Malloc1Size; - // Free memory should be less than assumed for - // single allocation smaller than min allocation chunk - REQUIRE(freeMemRet <= assumedFreeMem); - - HIP_CHECK(hipFree(A_mem)); -} TEST_CASE("Unit_hipMemGetInfo_DifferentMallocMultiSmall") { size_t freeMemInit; @@ -175,13 +194,52 @@ TEST_CASE("Unit_hipMemGetInfo_DifferentMallocMultiSmall") { auto assumedFreeMem = freeMemInit - (MallocSize * 2); + // freeMemRet should be FreeMem - (1 * MinAlloc) + // instead of FreeMem - (MinAlloc * 2) + // since MinAlloc > MallocSize*2 + REQUIRE(freeMemRet < assumedFreeMem); + fixAllocSize(MallocSize); + assumedFreeMem = freeMemInit - (MallocSize * 2); + // Ensure memory allocated is less than 2 * minimum allocation + REQUIRE(freeMemRet > assumedFreeMem); - // Confirm mem alocation results + // confirms that allocated memory is at least equal to Min Allocation + assumedFreeMem = freeMemInit - MinAlloc::Get(); REQUIRE(freeMemRet <= assumedFreeMem); HIP_CHECK(hipFree(A_mem)); HIP_CHECK(hipFree(B_mem)); } +TEST_CASE("Unit_hipMemGetInfo_DifferentMallocNotDiv") { + size_t freeMemInit; + size_t totalMemInit; + HIP_CHECK(hipMemGetInfo(&freeMemInit, &totalMemInit)); + + unsigned int* A_mem{nullptr}; + size_t freeMemRet; + size_t totalMemRet; + // Allocate memory that is just a bit larger than the min allocation + // Expected behaviour is to allocate 2x min allocation size + size_t MallocSize = MinAlloc::Get() + 1; + + HIP_CHECK(hipMalloc(reinterpret_cast(&A_mem), MallocSize)); + + HIP_CHECK(hipMemGetInfo(&freeMemRet, &totalMemRet)); + MEMINFO(totalMemRet, freeMemInit, freeMemRet, MallocSize); + + + auto freeMemExpected = freeMemInit - MallocSize; + // Free Memory after allocation should be less than + // expected free memory + REQUIRE(freeMemRet < freeMemExpected); + // confirms that allocated memory is at least 2 x Min Allocaton + fixAllocSize(MallocSize); + freeMemExpected = freeMemInit - MallocSize; + REQUIRE(freeMemRet <= freeMemExpected); + HIP_CHECK(hipFree(A_mem)); +} + + TEMPLATE_TEST_CASE("Unit_hipMemGetInfo_MallocArray", "", int, int4, char) { // get initial mem data size_t freeMemInit; @@ -209,6 +267,7 @@ TEMPLATE_TEST_CASE("Unit_hipMemGetInfo_MallocArray", "", int, int4, char) { size_t usedMem = bytesPerItem * extent.width * (extent.height != 0 ? extent.height : 1); // ensure we allocate at least the min allocation for the array + fixAllocSize(usedMem); MEMINFO(totalMemRet, freeMemInit, freeMemRet, usedMem); size_t assumedFreeMem = freeMemInit - usedMem; @@ -227,9 +286,9 @@ TEST_CASE("Unit_hipMemGetInfo_Malloc3D") { // Allocate 3D object hipExtent extent{}; // extent is given in bytes for with - extent.width = GENERATE(32, 128, 256); - extent.height = GENERATE(32, 128, 256); - extent.depth = GENERATE(32, 128, 256); + extent.width = GENERATE(32, 128, 256, 512); + extent.height = GENERATE(32, 128, 256, 512); + extent.depth = GENERATE(32, 128, 256, 512); hipPitchedPtr A_mem{}; HIP_CHECK(hipMalloc3D(&A_mem, extent)); @@ -240,6 +299,7 @@ TEST_CASE("Unit_hipMemGetInfo_Malloc3D") { // Verify result size_t mallocSize = A_mem.pitch * extent.height * extent.depth; + fixAllocSize(mallocSize); size_t assumedFreeMem = freeMemInit - mallocSize; MEMINFO(totalMemRet, freeMemInit, freeMemRet, mallocSize); @@ -313,6 +373,9 @@ TEMPLATE_TEST_CASE("Unit_hipMemGetInfo_Malloc3DArray", "", char, int, int4) { REQUIRE(mallocSize <= static_cast(MinAlloc::Get())); } else { + // account for min allocation + fixAllocSize(mallocSize); + MEMINFO(totalMemRet, freeMemInit, freeMemRet, mallocSize); size_t assumedFreeMem = freeMemInit - mallocSize; REQUIRE(freeMemRet <= assumedFreeMem); @@ -332,13 +395,15 @@ TEST_CASE("Unit_hipMemGetInfo_ParaLarge") { auto Malloc1Size = freeMemInit >> 1; // if the allocation is not divisible by the MinAllocation // take into account and add padding + fixAllocSize(Malloc1Size); std::thread t1( - [&] { HIP_CHECK_THREAD(hipMalloc(reinterpret_cast(&A_mem), Malloc1Size)); }); + [&]() { HIP_CHECK_THREAD(hipMalloc(reinterpret_cast(&A_mem), Malloc1Size)); }); // allocate an extra quarter of free mem auto Malloc2Size = Malloc1Size >> 1; + fixAllocSize(Malloc2Size); std::thread t2( - [&] { HIP_CHECK_THREAD(hipMalloc(reinterpret_cast(&B_mem), Malloc2Size)); }); + [&]() { HIP_CHECK_THREAD(hipMalloc(reinterpret_cast(&B_mem), Malloc2Size)); }); t1.join(); t2.join(); @@ -356,16 +421,18 @@ TEST_CASE("Unit_hipMemGetInfo_ParaLarge") { HIP_CHECK(hipFree(B_mem)); } +#endif + TEST_CASE("Unit_hipMemGetInfo_ParaSmall") { size_t freeMemInit; size_t totalMemInit; HIP_CHECK(hipMemGetInfo(&freeMemInit, &totalMemInit)); unsigned int* A_mem{nullptr}; // allocate smaller chunk than minimum - size_t Malloc1Size = 1; + size_t Malloc1Size = 2; std::thread t1( - [&] { HIP_CHECK_THREAD(hipMalloc(reinterpret_cast(&A_mem), Malloc1Size)) }); + [&]() { HIP_CHECK_THREAD(hipMalloc(reinterpret_cast(&A_mem), Malloc1Size)) }); t1.join(); HIP_CHECK_THREAD_FINALIZE(); size_t freeMemRet; @@ -377,13 +444,101 @@ TEST_CASE("Unit_hipMemGetInfo_ParaSmall") { auto assumedFreeMem = freeMemInit - Malloc1Size; // Free memory should be less than assumed for // single allocation smaller than min allocation chunk + REQUIRE(freeMemRet < assumedFreeMem); + // confirms that allocated memory is at least equal to smallest allocation allowed + assumedFreeMem = freeMemInit - MinAlloc::Get(); + REQUIRE(freeMemRet <= assumedFreeMem); + + HIP_CHECK(hipFree(A_mem)); + + // allocate smallest chunck of memory + std::thread t2( + [&]() { HIP_CHECK_THREAD(hipMalloc(reinterpret_cast(&A_mem), MinAlloc::Get())); }); + t2.join(); + HIP_CHECK_THREAD_FINALIZE(); + + HIP_CHECK(hipMemGetInfo(&freeMemRet, &totalMemRet)); + + MEMINFO(totalMemRet, freeMemInit, freeMemRet, MinAlloc::Get()); + + assumedFreeMem = freeMemInit - MinAlloc::Get(); REQUIRE(freeMemRet <= assumedFreeMem); HIP_CHECK(hipFree(A_mem)); } +TEST_CASE("Unit_hipMemGetInfo_ParaNonDiv") { + size_t freeMemInit; + size_t totalMemInit; + HIP_CHECK(hipMemGetInfo(&freeMemInit, &totalMemInit)); + unsigned int* A_mem{nullptr}; + + // Allocate memory that is just 1 byte larger than the min allocation + // Expected behaviour is to allocate 2x min allocation size + size_t Malloc1Size = MinAlloc::Get() + 1; + + std::thread t1( + [&]() { HIP_CHECK_THREAD(hipMalloc(reinterpret_cast(&A_mem), Malloc1Size)); }); + t1.join(); + HIP_CHECK_THREAD_FINALIZE(); + + size_t freeMemRet; + size_t totalMemRet; + HIP_CHECK(hipMemGetInfo(&freeMemRet, &totalMemRet)); + MEMINFO(totalMemRet, freeMemInit, freeMemRet, Malloc1Size); + + + auto allocSize = freeMemInit - Malloc1Size; + // should not be equal + REQUIRE(freeMemRet != allocSize); + // confirms that allocated memory is equal to 2 x Min Allocaton + allocSize = MinAlloc::Get() * 2; + auto assumedAllocSize = freeMemInit - allocSize; + REQUIRE(freeMemRet <= assumedAllocSize); + HIP_CHECK(hipFree(A_mem)); +} + +TEST_CASE("Unit_hipMemGetInfo_ParaMultiSmall") { + size_t freeMemInit; + size_t totalMemInit; + HIP_CHECK(hipMemGetInfo(&freeMemInit, &totalMemInit)); + unsigned int* A_mem{nullptr}; + unsigned int* B_mem{nullptr}; + + // Allocate memory that is a quarter of the min allocation + // Expected behaviour is to reuse the min allocation memory + size_t MallocSize = MinAlloc::Get() >> 2; + + std::thread t1( + [&]() { HIP_CHECK_THREAD(hipMalloc(reinterpret_cast(&A_mem), MallocSize)); }); + std::thread t2( + [&]() { HIP_CHECK_THREAD(hipMalloc(reinterpret_cast(&B_mem), MallocSize)); }); + + t1.join(); + t2.join(); + HIP_CHECK_THREAD_FINALIZE(); + + size_t freeMemRet; + size_t totalMemRet; + HIP_CHECK(hipMemGetInfo(&freeMemRet, &totalMemRet)); + MEMINFO(totalMemRet, freeMemInit, freeMemRet, MallocSize * 2); + + auto assumedFreeMem = freeMemInit - MallocSize * 2; + // freeMemRet should be less than assumedFreeMem + REQUIRE(freeMemRet < assumedFreeMem); + // confirms that allocated memory is equal to Min Allocation + assumedFreeMem = freeMemInit - MinAlloc::Get(); + REQUIRE(freeMemRet <= assumedFreeMem); + HIP_CHECK(hipFree(A_mem)); + HIP_CHECK(hipFree(B_mem)); +} + TEST_CASE("Unit_hipMemGetInfo_Negative") { +#if HT_AMD + HipTest::HIP_SKIP_TEST(" EXSWCPHIPT-61"); + return; +#endif size_t freeMemInit; size_t totalMemInit; HIP_CHECK(hipMemGetInfo(&freeMemInit, &totalMemInit)); @@ -414,10 +569,6 @@ TEST_CASE("Unit_hipMemGetInfo_Negative") { HIP_CHECK(hipMemGetInfo(&freeMemRet, totalMemRet)); } SECTION("Nullptr as both params passed to hipMemGetInfo") { -#if HT_AMD - HipTest::HIP_SKIP_TEST("EXSWCPHIPT-135"); - return; -#endif size_t* freeMemRet = nullptr; size_t* totalMemRet = nullptr; HIP_CHECK(hipMalloc(reinterpret_cast(&A_mem), MallocSize)); diff --git a/tests/catch/unit/memory/hipMemcpySync.cc b/tests/catch/unit/memory/hipMemcpySync.cc new file mode 100644 index 0000000000..fbb165d892 --- /dev/null +++ b/tests/catch/unit/memory/hipMemcpySync.cc @@ -0,0 +1,227 @@ +/* +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of intge, 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 "MemUtils.hh" + +/* + * These testcases verify that synchronization behaviour for memcpy functions with respect to + * the host. + */ + +using namespace mem_utils; + +// value used for memset operations +constexpr int testValue = 0x11; + + +/* + * Set of helper functions handling the different cases for memcpy + */ + +static inline hipMemcpyKind getMemcpyType(allocType type, bool fromHost) { + if (fromHost) { + switch (type) { + case allocType::deviceMalloc: + return hipMemcpyHostToDevice; + break; + case allocType::devRegistered: + return hipMemcpyHostToDevice; + break; + default: // host + return hipMemcpyHostToHost; + break; + } + } else { + switch (type) { + case allocType::deviceMalloc: + return hipMemcpyDeviceToDevice; + break; + case allocType::devRegistered: + return hipMemcpyDeviceToDevice; + break; + default: // host + return hipMemcpyDeviceToHost; + break; + } + } +} + +static inline void memcpyCheck(allocType type, memType memType, char* aPtr, MultiDData& data, + char* fillerData, bool async, hipStream_t stream, bool fromHost) { + auto cpyType = getMemcpyType(type, fromHost); + auto sizeInBytes = data.pitch * data.getH() * data.getD() * sizeof(char); + switch (memType) { + case memType::hipMem: + if (async) { + HIP_CHECK(hipMemcpyAsync(aPtr + data.offset, fillerData, sizeInBytes, cpyType, stream)); + } else { + HIP_CHECK(hipMemcpy(aPtr + data.offset, fillerData, sizeInBytes, cpyType)); + } + break; + case memType::hipMem2D: + if (async) { + HIP_CHECK(hipMemcpy2DAsync(aPtr + data.offset, data.pitch, fillerData, sizeInBytes, + data.width, data.getH(), cpyType, stream)); + } else { + HIP_CHECK(hipMemcpy2D(aPtr + data.offset, data.pitch, fillerData, sizeInBytes, data.width, + data.getH(), cpyType)); + } + break; + case memType::hipMem3D: { + hipMemcpy3DParms params{}; + params.kind = cpyType; + params.srcPos = make_hipPos(0, 0, 0); + params.dstPos = make_hipPos(data.offset, data.offset, data.offset); + params.srcPtr = make_hipPitchedPtr(fillerData, data.width, data.width, data.getH()); + params.dstPtr = make_hipPitchedPtr(aPtr, data.pitch, data.width, data.getH()); + hipExtent extent; + extent.width = data.width * sizeof(char); + extent.height = data.getH(); + extent.depth = data.getD(); + + params.extent = extent; + if (async) { + HIP_CHECK(hipMemcpy3DAsync(¶ms, stream)); + } else { + HIP_CHECK(hipMemcpy3D(¶ms)); + } + break; + } + default: + break; + } +} + +static inline char* createFillerData(size_t count, size_t value, bool fromHost) { + if (fromHost) { + char* fillerData = new char[count]; + std::fill(fillerData, fillerData + count, value); + return fillerData; + } else { + char* fillerData; + HIP_CHECK(hipMalloc(&fillerData, count * sizeof(char))); + HIP_CHECK(hipMemset(fillerData, value, count * sizeof(char))); + return fillerData; + } +} + +static void checkForSync(hipStream_t stream, bool async, allocType type, bool fromHost) { + if (fromHost) { + if (type == allocType::deviceMalloc) { + HIP_CHECK_ERROR(hipStreamQuery(stream), hipErrorNotReady); + } else { + HIP_CHECK(hipStreamQuery(stream)); + } + } else { + if (type != allocType::deviceMalloc && !async) { + HIP_CHECK(hipStreamQuery(stream)); + } else { + HIP_CHECK_ERROR(hipStreamQuery(stream), hipErrorNotReady); + } + } +} + + +// Helper function to run tests for hipMemset allocation types +static void runMemcpyTests(hipStream_t stream, bool async, allocType type, memType memType, + MultiDData data) { + bool fromHost = GENERATE(true, false); + + std::pair aPtr = initMemory(type, memType, data); + size_t sizeInBytes = data.getCount(); + + // filler data for device memory created beforehand as it uses memset + // which might interfere with synchronization testing + auto fillerData = createFillerData(sizeInBytes, testValue, fromHost); + CAPTURE(type, memType, data.width, data.height, data.depth, stream, async, fromHost, sizeInBytes); + + using namespace std::chrono_literals; + const std::chrono::duration delay = 100ms; + HipTest::runKernelForDuration(delay, stream); + + memcpyCheck(type, memType, aPtr.first, data, fillerData, async, stream, fromHost); + checkForSync(stream, async, type, fromHost); + // verify + HIP_CHECK(hipStreamSynchronize(stream)); + verifyData(aPtr.first, testValue, data, type, memType); + if (type == allocType::devRegistered) { + freeStuff(aPtr.second, type); + } else { + freeStuff(aPtr.first, type); + } + if (fromHost) { + delete[] fillerData; + } else { + HIP_CHECK(hipFree(fillerData)); + } +} + +#if HT_AMD /* Disabled because frequency based wait is timing out on nvidia platforms */ + +TEST_CASE("Unit_hipMemcpySync") { +#if HT_AMD // To be removed when EXSWCPHIPT-127 is fixed + HipTest::HIP_SKIP_TEST("EXSWCPHIPT-127 - Sync behaviour differs on AMD and Nvidia"); + return; +#endif + allocType type = GENERATE(allocType::deviceMalloc, allocType::hostMalloc, allocType::hostRegisted, + allocType::devRegistered); + memType memcpy_type = memType::hipMem; + MultiDData data; + data.width = 1; + + doMemTest(runMemcpyTests, type, memcpy_type, data); // Uses long running kernel +} + +TEST_CASE("Unit_hipMemcpy2DSync") { +#if HT_AMD + HipTest::HIP_SKIP_TEST("EXSWCPHIPT-127 - Sync behaviour differs on AMD and Nvidia"); + return; +#endif + allocType mallocType = GENERATE(allocType::deviceMalloc, allocType::hostMalloc, + allocType::hostRegisted, allocType::devRegistered); + + memType memcpy_type = memType::hipMem2D; + MultiDData data; + data.width = 1; + data.height = 1; + + doMemTest(runMemcpyTests, mallocType, memcpy_type, data); +} + +TEST_CASE("Unit_hipMemcpy3DSync") { +#if HT_AMD + HipTest::HIP_SKIP_TEST("EXSWCPHIPT-127 - Sync behaviour differs on AMD and Nvidia"); + return; +#endif + allocType mallocType = GENERATE(allocType::deviceMalloc, allocType::hostMalloc, + allocType::hostRegisted, allocType::devRegistered); + + memType memcpy_type = memType::hipMem3D; + MultiDData data; + data.width = 1; + data.height = 1; + data.depth = 1; + + doMemTest(runMemcpyTests, mallocType, memcpy_type, data); +} + +#endif diff --git a/tests/catch/unit/memory/hipMemsetAsync.cc b/tests/catch/unit/memory/hipMemsetAsync.cc new file mode 100644 index 0000000000..b0752e3694 --- /dev/null +++ b/tests/catch/unit/memory/hipMemsetAsync.cc @@ -0,0 +1,173 @@ +/* +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of intge, 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 "MemUtils.hh" + +/* + * This testcase verifies that asynchronous memset functions are asynchronous with respect to the + * host except when the target is pinned host memory or a Unified Memory region + */ + +constexpr int testValue1 = 97; +constexpr int testValue2 = 98; + + +using namespace mem_utils; + +// Helper function to run tests for hipMemset allocation types +template +void runAsyncTests(hipStream_t stream, allocType type, memType memType, MultiDData data1, + MultiDData data2) { + std::pair aPtr{}; + MultiDData totalRange; + totalRange.width = data1.width + data2.width; + totalRange.height = data1.height + data2.height; + totalRange.depth = data1.depth + data2.depth; + aPtr = initMemory(type, memType, totalRange); + data1.pitch = totalRange.pitch; + data2.pitch = totalRange.pitch; + + memsetCheck(aPtr.first, testValue1, memType, data1, stream); + memsetCheck(aPtr.first, testValue2, memType, data2, stream); + + HIP_CHECK(hipStreamSynchronize(stream)); + verifyData(aPtr.first, testValue1, data1, type, memType); + verifyData(aPtr.first, testValue2, data2, type, memType); + + + if (type == allocType::devRegistered) { + freeStuff(aPtr.second, type); + } else { + freeStuff(aPtr.first, type); + } +} + +template +static void doMemsetTest(allocType mallocType, memType memset_type, MultiDData data1, + MultiDData data2) { + enum StreamType { NULLSTR, CREATEDSTR }; + auto streamType = GENERATE(NULLSTR, CREATEDSTR); + hipStream_t stream{nullptr}; + + if (streamType == CREATEDSTR) HIP_CHECK(hipStreamCreate(&stream)); + + runAsyncTests(stream, mallocType, memset_type, data1, data2); + + if (streamType == CREATEDSTR) HIP_CHECK(hipStreamDestroy(stream)); +} + +/* + * test 2 async hipMemset's on the same memory at different offsets + */ + +TEST_CASE("Unit_hipMemsetASyncMulti") { +#if HT_AMD + HipTest::HIP_SKIP_TEST("EXSWCPHIPT-127"); + return; +#endif + allocType mallocType = GENERATE(allocType::hostMalloc, allocType::deviceMalloc, + allocType::hostRegisted, allocType::devRegistered); + memType mem_type = memType::hipMemsetD8; + MultiDData data1; + data1.offset = 0; + data1.width = GENERATE(1, 256); + MultiDData data2; + data2.width = data1.width; + + data2.offset = data1.width; + doMemsetTest(mallocType, mem_type, data1, data2); +} + +/* + * test 2 async hipMemsetD[8,16,32]'s on the same memory at different offsets + */ +TEMPLATE_TEST_CASE("Unit_hipMemsetDASyncMulti", "", int8_t, int16_t, uint32_t) { +#if HT_AMD + HipTest::HIP_SKIP_TEST("EXSWCPHIPT-127"); + return; +#endif + allocType mallocType = GENERATE(allocType::hostRegisted, allocType::deviceMalloc, + allocType::hostMalloc, allocType::devRegistered); + memType memset_type; + MultiDData data1; + data1.offset = 0; + data1.width = GENERATE(1, 256); + MultiDData data2; + data2.width = data1.width; + data2.offset = data1.width; + + if (std::is_same::value) { + memset_type = memType::hipMemsetD8; + } else if (std::is_same::value) { + memset_type = memType::hipMemsetD16; + } else if (std::is_same::value) { + memset_type = memType::hipMemsetD32; + } + doMemsetTest(mallocType, memset_type, data1, data2); +} + +/* + * test 2 async hipMemset2D's on the same memory at different offsets + */ +TEST_CASE("Unit_hipMemset2DASyncMulti") { +#if HT_AMD + HipTest::HIP_SKIP_TEST("EXSWCPHIPT-127"); + return; +#endif + allocType mallocType = GENERATE(allocType::deviceMalloc, allocType::hostMalloc, + allocType::hostRegisted, allocType::devRegistered); + memType memset_type = memType::hipMem2D; + MultiDData data1; + data1.offset = 0; + data1.width = GENERATE(1, 256); + data1.height = data1.width; + MultiDData data2; + data2.width = data1.width; + data2.height = data1.height; + data2.offset = data1.width; + + doMemsetTest(mallocType, memset_type, data1, data2); +} +/* + * test 2 async hipMemset3D's on the same memory at different offsets + */ +TEST_CASE("Unit_hipMemset3DASyncMulti") { +#if HT_AMD + HipTest::HIP_SKIP_TEST("EXSWCPHIPT-127"); + return; +#endif + allocType mallocType = GENERATE(allocType::deviceMalloc, allocType::hostMalloc, + allocType::hostRegisted, allocType::devRegistered); + memType memset_type = memType::hipMem3D; + MultiDData data1; + data1.offset = 0; + data1.width = GENERATE(1, 256); + data1.height = data1.width; + data1.depth = data1.width; + MultiDData data2; + data2.width = data1.width; + data2.height = data1.width; + data2.depth = data1.width; + data2.offset = data1.width; + + doMemsetTest(mallocType, memset_type, data1, data2); +} diff --git a/tests/catch/unit/memory/hipMemsetSync.cc b/tests/catch/unit/memory/hipMemsetSync.cc new file mode 100644 index 0000000000..6770a6e16c --- /dev/null +++ b/tests/catch/unit/memory/hipMemsetSync.cc @@ -0,0 +1,505 @@ +/* +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of intge, 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 +/* + * These testcases verify that synchronous memset functions are asynchronous with respect to the + * host except when the target is pinned host memory or a Unified Memory region + */ + +// value used for memset operations +constexpr int testValue = 0x11; + +enum class allocType { deviceMalloc, hostMalloc, hostRegisted, devRegistered }; +enum class memSetType { + hipMemset, + hipMemsetD8, + hipMemsetD16, + hipMemsetD32, + hipMemset2D, + hipMemset3D +}; + +// helper struct containing vars needed for 2D and 3D memset Testing +struct MultiDData { + size_t width{}; + // set to 0 for 1D + size_t height{}; + // set to 0 for 2D + size_t depth{}; + size_t pitch{}; +}; + +// set of helper functions to tidy the nested switch statements +template +static std::pair deviceMallocHelper(memSetType memType, size_t dataW, size_t dataH, size_t dataD, + size_t& dataPitch) { + size_t elementSize = sizeof(T); + size_t sizeInBytes = elementSize * dataW * dataH * dataD; + T* aPtr{}; + switch (memType) { + case memSetType::hipMemset3D: { + hipPitchedPtr pitchedAPtr{}; + hipExtent extent; + extent.width = dataW * elementSize; + extent.height = dataH; + extent.depth = dataD; + + pitchedAPtr = + make_hipPitchedPtr(aPtr, extent.width, extent.width / elementSize, extent.height); + HIP_CHECK(hipMalloc3D(&pitchedAPtr, extent)); + aPtr = reinterpret_cast(pitchedAPtr.ptr); + dataPitch = pitchedAPtr.pitch; + break; + } + + case memSetType::hipMemset2D: + HIP_CHECK( + hipMallocPitch(reinterpret_cast(&aPtr), &dataPitch, dataW * elementSize, dataH)); + + dataPitch = dataW * elementSize; + break; + + default: + HIP_CHECK(hipMalloc(&aPtr, sizeInBytes)); + dataPitch = dataW * elementSize; + break; + } + return std::make_pair(aPtr, nullptr); +} + +template +static std::pair hostMallocHelper(size_t dataW, size_t dataH, size_t dataD, size_t& dataPitch) { + size_t elementSize = sizeof(T); + size_t sizeInBytes = elementSize * dataW * dataH * dataD; + T* aPtr; + + HIP_CHECK(hipHostMalloc(&aPtr, sizeInBytes)); + dataPitch = dataW * elementSize; + + return std::make_pair(aPtr, nullptr); +} + +template +static std::pair hostRegisteredHelper(size_t dataW, size_t dataH, size_t dataD, size_t& dataPitch) { + size_t elementSize = sizeof(T); + size_t sizeInBytes = elementSize * dataW * dataH * dataD; + T* aPtr = new T[dataW * dataH * dataD]; + + HIP_CHECK(hipHostRegister(aPtr, sizeInBytes, hipHostRegisterDefault)); + + dataPitch = dataW * elementSize; + return std::make_pair(aPtr, nullptr); +} + +template +static std::pair devRegisteredHelper(size_t dataW, size_t dataH, size_t dataD, + size_t& dataPitch) { + size_t elementSize = sizeof(T); + size_t sizeInBytes = elementSize * dataW * dataH * dataD; + T* aPtr = new T[dataW * dataH * dataD]; + T* retPtr; + + HIP_CHECK(hipHostRegister(aPtr, sizeInBytes, hipHostRegisterDefault)); + HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&retPtr), aPtr, 0)); + + dataPitch = dataW * elementSize; + // keep the address of the host memory + return std::make_pair(retPtr, aPtr); +} + +// helper function to allocate memory and set it to a value. +// retunr a pair of pointers due to the device registered allocation case, we need to keep track of +// the pointer to host memory to be able to unregister and free it +template +static std::pair initMemory(allocType type, memSetType memType, MultiDData& data) { + size_t dataH = data.height == 0 ? 1 : data.height; + size_t dataD = data.depth == 0 ? 1 : data.depth; + std::pair retPtr{}; + // check different types of allocation + switch (type) { + case allocType::deviceMalloc: + retPtr = deviceMallocHelper(memType, data.width, dataH, dataD, data.pitch); + break; + + case allocType::hostMalloc: + retPtr = hostMallocHelper(data.width, dataH, dataD, data.pitch); + break; + + case allocType::hostRegisted: + retPtr = hostRegisteredHelper(data.width, dataH, dataD, data.pitch); + break; + + case allocType::devRegistered: + retPtr = devRegisteredHelper(data.width, dataH, dataD, data.pitch); + break; + + default: + REQUIRE(false); + break; + } + return retPtr; +} + +// set of helper functions to tidy the nested switch statements +template +static void deviceMallocCopy(memSetType memType, T* aPtr, T* hostMem, size_t dataW, size_t dataH, + size_t dataD, size_t& dataPitch) { + size_t elementSize = sizeof(T); + size_t sizeInBytes = elementSize * dataW * dataH * dataD; + switch (memType) { + case memSetType::hipMemset3D: { + hipMemcpy3DParms params{}; + params.kind = hipMemcpyDeviceToHost; + params.srcPos = make_hipPos(0, 0, 0); + params.srcPtr = make_hipPitchedPtr(aPtr, dataPitch, dataW, dataH); + params.dstPos = make_hipPos(0, 0, 0); + params.dstPtr = make_hipPitchedPtr(hostMem, dataPitch, dataW, dataH); + + hipExtent extent; + extent.width = dataPitch; + extent.height = dataH; + extent.depth = dataD; + + params.extent = extent; + + HIP_CHECK(hipMemcpy3D(¶ms)); + break; + } + + case memSetType::hipMemset2D: + HIP_CHECK(hipMemcpy2D(hostMem, dataW * elementSize, aPtr, dataPitch, dataW, dataH, + hipMemcpyDeviceToHost)); + break; + + default: + HIP_CHECK(hipMemcpy(hostMem, aPtr, sizeInBytes, hipMemcpyDeviceToHost)); + break; + } +} + +template +static void hostCopy(memSetType memType, T* aPtr, T* hostMem, size_t dataW, size_t dataH, + size_t dataD, size_t& dataPitch) { + size_t elementSize = sizeof(T); + size_t sizeInBytes = elementSize * dataW * dataH * dataD; + hipMemcpy3DParms params{}; + switch (memType) { + case memSetType::hipMemset3D: { + params.kind = hipMemcpyHostToHost; + params.srcPos = make_hipPos(0, 0, 0); + params.dstPos = make_hipPos(0, 0, 0); + params.srcPtr = make_hipPitchedPtr(aPtr, dataPitch, dataW, dataH); + params.dstPtr = make_hipPitchedPtr(hostMem, dataW, dataW, dataH); + + hipExtent extent; + extent.width = dataW; + extent.height = dataH; + extent.depth = dataD; + + params.extent = extent; + + HIP_CHECK(hipMemcpy3D(¶ms)); + break; + } + + case memSetType::hipMemset2D: + HIP_CHECK(hipMemcpy2D(hostMem, dataW * elementSize, aPtr, dataPitch, dataW, dataH, + hipMemcpyHostToHost)); + break; + + default: + HIP_CHECK(hipMemcpy(hostMem, aPtr, sizeInBytes, hipMemcpyHostToHost)); + break; + } +} + +template +static void devRegisteredCopy(memSetType memType, T* aPtr, T* hostMem, size_t dataW, size_t dataH, + size_t dataD, size_t& dataPitch) { + size_t elementSize = sizeof(T); + + switch (memType) { + case memSetType::hipMemset3D: { + hipMemcpy3DParms params{}; + params.kind = hipMemcpyHostToHost; + params.srcPos = make_hipPos(0, 0, 0); + params.dstPos = make_hipPos(0, 0, 0); + params.srcPtr = make_hipPitchedPtr(aPtr, dataPitch, dataW, dataH); + params.dstPtr = make_hipPitchedPtr(hostMem, dataW, dataW, dataH); + + hipExtent extent; + extent.width = dataW; + extent.height = dataH; + extent.depth = dataD; + + params.extent = extent; + + HIP_CHECK(hipMemcpy3D(¶ms)); + break; + } + + case memSetType::hipMemset2D: + HIP_CHECK(hipMemcpy2D(hostMem, dataW * elementSize, aPtr, dataPitch, dataW, dataH, + hipMemcpyDeviceToHost)); + break; + + default: { + size_t sizeInBytes = elementSize * dataW * dataH * dataD; + HIP_CHECK(hipMemcpy(hostMem, aPtr, sizeInBytes, hipMemcpyDeviceToHost)); + break; + } + } +} + +// Copies device data to host and checks that each element is equal to the +// specified value +template +void verifyData(T* aPtr, size_t value, MultiDData& data, allocType type, memSetType memType) { + auto dataH = data.height == 0 ? 1 : data.height; + auto dataD = data.depth == 0 ? 1 : data.depth; + std::unique_ptr hostPtr = std::make_unique(data.pitch * dataH * dataD / sizeof(T)); + switch (type) { + case allocType::deviceMalloc: + deviceMallocCopy(memType, aPtr, hostPtr.get(), data.width, dataH, dataD, data.pitch); + break; + case allocType::devRegistered: + devRegisteredCopy(memType, aPtr, hostPtr.get(), data.width, dataH, dataD, data.pitch); + break; + default: // host allocated or host registered memory + hostCopy(memType, aPtr, hostPtr.get(), data.width, dataH, dataD, data.pitch); + break; + } + + + size_t idx; + bool allMatch = true; + + for (size_t k = 0; k < dataD; k++) { + for (size_t j = 0; j < dataH; j++) { + for (size_t i = 0; i < data.width; i++) { + idx = data.pitch * dataH * k + data.pitch * j + i; + allMatch = allMatch && static_cast(hostPtr.get()[idx]) == value; + if (!allMatch) REQUIRE(false); + } + } + } +} + +// macro to allow reuse of functions for testing versions of hipMemset +template +void memsetCheck(T* aPtr, size_t value, memSetType memsetType, MultiDData& data, bool async = false, + hipStream_t stream = nullptr) { + size_t dataW = data.width; + size_t dataH = data.height == 0 ? 1 : data.height; + size_t dataD = data.depth == 0 ? 1 : data.depth; + size_t count = dataW * dataH * dataD; + + switch (memsetType) { + case memSetType::hipMemset: + if (async) { + HIP_CHECK(hipMemsetAsync(aPtr, value, count, stream)); + } else { + HIP_CHECK(hipMemset(aPtr, value, count)); + } + break; + + case memSetType::hipMemsetD8: + if (async) { + HIP_CHECK(hipMemsetD8Async(reinterpret_cast(aPtr), value, count, stream)); + } else { + HIP_CHECK(hipMemsetD8(reinterpret_cast(aPtr), value, count)); + } + + break; + + case memSetType::hipMemsetD16: + if (async) { + HIP_CHECK(hipMemsetD16Async(reinterpret_cast(aPtr), value, count, stream)); + } else { + HIP_CHECK(hipMemsetD16(reinterpret_cast(aPtr), value, count)); + } + break; + + case memSetType::hipMemsetD32: + if (async) { + HIP_CHECK(hipMemsetD32Async(reinterpret_cast(aPtr), value, count, stream)); + } else { + HIP_CHECK(hipMemsetD32(reinterpret_cast(aPtr), value, count)); + } + break; + + case memSetType::hipMemset2D: + if (async) { + HIP_CHECK(hipMemset2DAsync(aPtr, data.pitch, value, data.width, data.height, stream)); + } else { + HIP_CHECK(hipMemset2D(aPtr, data.pitch, value, data.width, data.height)); + } + break; + + case memSetType::hipMemset3D: + hipExtent extent; + extent.width = data.width; + extent.height = data.height; + extent.depth = data.depth; + if (async) { + HIP_CHECK(hipMemset3DAsync(make_hipPitchedPtr(aPtr, data.pitch, data.width, data.height), + value, extent, stream)); + } else { + HIP_CHECK(hipMemset3D(make_hipPitchedPtr(aPtr, data.pitch, data.width, data.height), value, + extent)); + } + break; + + default: + REQUIRE(false); + break; + } +} + +template void freeStuff(T* aPtr, allocType type) { + switch (type) { + case allocType::deviceMalloc: + hipFree(aPtr); + break; + case allocType::hostMalloc: + hipHostFree(aPtr); + break; + case allocType::hostRegisted: + HIP_CHECK(hipHostUnregister(aPtr)); + delete[] aPtr; + break; + case allocType::devRegistered: + HIP_CHECK(hipHostUnregister(aPtr)); + delete[] aPtr; + break; + default: + REQUIRE(false); + break; + } +} + +// Helper function to run tests for hipMemset allocation types +template +void runTests(allocType type, memSetType memsetType, MultiDData data, hipStream_t stream) { + bool async = GENERATE(true, false); + CAPTURE(type, memsetType, data.width, data.height, data.depth, stream, async); + std::pair aPtr = initMemory(type, memsetType, data); + using namespace std::chrono_literals; + const std::chrono::duration delay = 100ms; + HipTest::runKernelForDuration(delay, stream); + memsetCheck(aPtr.first, testValue, memsetType, data, async, stream); + + if (async || type == allocType::deviceMalloc) { + HIP_CHECK_ERROR(hipStreamQuery(stream), hipErrorNotReady); + } else { + HIP_CHECK(hipStreamQuery(stream)); + } + + HIP_CHECK(hipStreamSynchronize(stream)); + verifyData(aPtr.first, testValue, data, type, memsetType); + + if (type == allocType::devRegistered) { + freeStuff(aPtr.second, type); + } else { + freeStuff(aPtr.first, type); + } +} + +template +static void doMemsetTest(allocType mallocType, memSetType memset_type, MultiDData data) { + enum StreamType { NULLSTR, CREATEDSTR }; + auto streamType = GENERATE(NULLSTR, CREATEDSTR); + hipStream_t stream{nullptr}; + + if (streamType == CREATEDSTR) HIP_CHECK(hipStreamCreate(&stream)); + + runTests(mallocType, memset_type, data, stream); + + if (streamType == CREATEDSTR) HIP_CHECK(hipStreamDestroy(stream)); +} + +TEST_CASE("Unit_hipMemsetSync") { +#if HT_AMD || HT_NVIDIA + HipTest::HIP_SKIP_TEST("EXSWCPHIPT-86"); + return; +#endif + allocType type = GENERATE(allocType::deviceMalloc, allocType::hostMalloc, allocType::hostRegisted, + allocType::devRegistered); + memSetType memset_type = memSetType::hipMemset; + MultiDData data; + data.width = GENERATE(1, 1024); + doMemsetTest(type, memset_type, data); +} + +TEMPLATE_TEST_CASE("Unit_hipMemsetDSync", "", int8_t, int16_t, uint32_t) { +#if HT_AMD || HT_NVIDIA + HipTest::HIP_SKIP_TEST("EXSWCPHIPT-86"); + return; +#endif + allocType mallocType = GENERATE(allocType::hostRegisted, allocType::deviceMalloc, + allocType::hostMalloc, allocType::devRegistered); + memSetType memset_type; + MultiDData data; + data.width = GENERATE(1, 1024); + + if (std::is_same::value) { + memset_type = memSetType::hipMemsetD8; + } else if (std::is_same::value) { + memset_type = memSetType::hipMemsetD16; + } else if (std::is_same::value) { + memset_type = memSetType::hipMemsetD32; + } + + doMemsetTest(mallocType, memset_type, data); +} + +TEST_CASE("Unit_hipMemset2DSync") { +#if HT_AMD || HT_NVIDIA + HipTest::HIP_SKIP_TEST("EXSWCPHIPT-86"); + return; +#endif + allocType mallocType = GENERATE(allocType::deviceMalloc, allocType::hostMalloc, + allocType::hostRegisted, allocType::devRegistered); + memSetType memset_type = memSetType::hipMemset2D; + MultiDData data; + data.width = GENERATE(1, 1024); + data.height = GENERATE(1, 1024); + + doMemsetTest(mallocType, memset_type, data); +} + +TEST_CASE("Unit_hipMemset3DSync") { +#if HT_AMD || HT_NVIDIA + HipTest::HIP_SKIP_TEST("EXSWCPHIPT-86"); + return; +#endif + allocType mallocType = GENERATE(allocType::deviceMalloc, allocType::hostMalloc, + allocType::hostRegisted, allocType::devRegistered); + memSetType memset_type = memSetType::hipMemset3D; + MultiDData data; + data.width = GENERATE(1, 256); + data.height = GENERATE(1, 256); + data.depth = GENERATE(1, 256); + + doMemsetTest(mallocType, memset_type, data); +} diff --git a/tests/catch/unit/memory/hipPointerGetAttributes.cc b/tests/catch/unit/memory/hipPointerGetAttributes.cc index 1e02b696fc..8157c8811d 100644 --- a/tests/catch/unit/memory/hipPointerGetAttributes.cc +++ b/tests/catch/unit/memory/hipPointerGetAttributes.cc @@ -28,7 +28,7 @@ Following scenarios are verified for hipPointerGetAttributes API 4. Multi-threaded test with many simul allocs. */ -#include +#include #include #include #include @@ -37,22 +37,18 @@ size_t Nbytes = 0; constexpr size_t N{1000000}; - //================================================================================================= // Utility Functions: //================================================================================================= -bool operator==(const hipPointerAttribute_t& lhs, - const hipPointerAttribute_t& rhs) { - return ((lhs.hostPointer == rhs.hostPointer) && - (lhs.devicePointer == rhs.devicePointer) && - (lhs.memoryType == rhs.memoryType) && (lhs.device == rhs.device) && - (lhs.allocationFlags == rhs.allocationFlags)); +bool operator==(const hipPointerAttribute_t& lhs, const hipPointerAttribute_t& rhs) { + return ((lhs.hostPointer == rhs.hostPointer) && (lhs.devicePointer == rhs.devicePointer) && + (lhs.memoryType == rhs.memoryType) && (lhs.device == rhs.device) && + (lhs.allocationFlags == rhs.allocationFlags)); } -bool operator!=(const hipPointerAttribute_t& lhs, - const hipPointerAttribute_t& rhs) { +bool operator!=(const hipPointerAttribute_t& lhs, const hipPointerAttribute_t& rhs) { return !(lhs == rhs); } @@ -70,53 +66,50 @@ const char* memoryTypeToString(hipMemoryType memoryType) { void resetAttribs(hipPointerAttribute_t* attribs) { - attribs->hostPointer = reinterpret_cast(-1); - attribs->devicePointer = reinterpret_cast(-1); - attribs->memoryType = hipMemoryTypeHost; - attribs->device = -2; - attribs->isManaged = -1; - attribs->allocationFlags = 0xffff; + attribs->hostPointer = reinterpret_cast(-1); + attribs->devicePointer = reinterpret_cast(-1); + attribs->memoryType = hipMemoryTypeHost; + attribs->device = -2; + attribs->isManaged = -1; + attribs->allocationFlags = 0xffff; } void printAttribs(const hipPointerAttribute_t* attribs) { printf( - "hostPointer:%p devicePointer:%p memType:%s deviceId:%d isManaged:%d " - "allocationFlags:%u\n", - attribs->hostPointer, attribs->devicePointer, - memoryTypeToString(attribs->memoryType), - attribs->device, attribs->isManaged, attribs->allocationFlags); + "hostPointer:%p devicePointer:%p memType:%s deviceId:%d isManaged:%d " + "allocationFlags:%u\n", + attribs->hostPointer, attribs->devicePointer, memoryTypeToString(attribs->memoryType), + attribs->device, attribs->isManaged, attribs->allocationFlags); } inline int zrand(int max) { return rand() % max; } - // Store the hipPointer attrib and some extra info // so can later compare the looked-up info against // the reference expectation struct SuperPointerAttribute { - void* _pointer; - size_t _sizeBytes; - hipPointerAttribute_t _attrib; + void* _pointer; + size_t _sizeBytes; + hipPointerAttribute_t _attrib; }; // Support function to check result against a reference: -void checkPointer(const SuperPointerAttribute& ref, int major, - int minor, void* pointer) { - hipPointerAttribute_t attribs; - resetAttribs(&attribs); +void checkPointer(const SuperPointerAttribute& ref, int major, int minor, void* pointer) { + hipPointerAttribute_t attribs; + resetAttribs(&attribs); - hipError_t e = hipPointerGetAttributes(&attribs, pointer); - if ((e != hipSuccess) || (attribs != ref._attrib)) { - HIP_CHECK(e); - REQUIRE(attribs != ref._attrib); - } else { - printf("#%4d.%d GOOD:%p getattr :: ", major, minor, pointer); - printAttribs(&attribs); - } + hipError_t e = hipPointerGetAttributes(&attribs, pointer); + if ((e != hipSuccess) || (attribs != ref._attrib)) { + HIP_CHECK(e); + REQUIRE(attribs != ref._attrib); + } else { + printf("#%4d.%d GOOD:%p getattr :: ", major, minor, pointer); + printAttribs(&attribs); + } } @@ -129,8 +122,7 @@ void checkPointer(const SuperPointerAttribute& ref, int major, // we do this in the testMultiThreaded_1 test. void clusterAllocs(int numAllocs, size_t minSize, size_t maxSize) { Nbytes = N * sizeof(char); - printf("clusterAllocs numAllocs=%d size=%lu..%lu\n", - numAllocs, minSize, maxSize); + printf("clusterAllocs numAllocs=%d size=%lu..%lu\n", numAllocs, minSize, maxSize); const int Max_Devices = 256; std::vector reference(numAllocs); @@ -157,18 +149,15 @@ void clusterAllocs(int numAllocs, size_t minSize, size_t maxSize) { void* ptr; if (isDevice) { - totalDeviceAllocated[reference[i]._attrib.device] += - reference[i]._sizeBytes; - HIP_CHECK(hipMalloc(reinterpret_cast(&ptr), - reference[i]._sizeBytes)); + totalDeviceAllocated[reference[i]._attrib.device] += reference[i]._sizeBytes; + HIP_CHECK(hipMalloc(reinterpret_cast(&ptr), reference[i]._sizeBytes)); reference[i]._attrib.memoryType = hipMemoryTypeDevice; reference[i]._attrib.devicePointer = ptr; reference[i]._attrib.hostPointer = NULL; reference[i]._attrib.allocationFlags = 0; } else { - HIP_CHECK(hipHostMalloc(reinterpret_cast(&ptr), - reference[i]._sizeBytes, - hipHostMallocDefault)); + HIP_CHECK(hipHostMalloc(reinterpret_cast(&ptr), reference[i]._sizeBytes, + hipHostMallocDefault)); reference[i]._attrib.memoryType = hipMemoryTypeHost; reference[i]._attrib.devicePointer = ptr; reference[i]._attrib.hostPointer = ptr; @@ -182,32 +171,29 @@ void clusterAllocs(int numAllocs, size_t minSize, size_t maxSize) { HIP_CHECK(hipSetDevice(i)); HIP_CHECK(hipMemGetInfo(&free, &total)); printf( - " device#%d: hipMemGetInfo: " - "free=%zu (%4.2fMB) totalDevice=%lu (%4.2fMB) total=%zu " - "(%4.2fMB)\n", - i, free, (free / 1024.0 / 1024.0), totalDeviceAllocated[i], - (totalDeviceAllocated[i]) / 1024.0 / 1024.0, total, - (total / 1024.0 / 1024.0)); + " device#%d: hipMemGetInfo: " + "free=%zu (%4.2fMB) totalDevice=%lu (%4.2fMB) total=%zu " + "(%4.2fMB)\n", + i, free, (free / 1024.0 / 1024.0), totalDeviceAllocated[i], + (totalDeviceAllocated[i]) / 1024.0 / 1024.0, total, (total / 1024.0 / 1024.0)); REQUIRE(free + totalDeviceAllocated[i] <= total); } // Now look up each pointer we inserted and verify we can find it: - char * ptr; + char* ptr; for (int i = 0; i < numAllocs; i++) { SuperPointerAttribute& ref = reference[i]; - ptr = static_cast(ref._pointer); + ptr = static_cast(ref._pointer); checkPointer(ref, i, 0, ref._pointer); - checkPointer(ref, i, 1, (ptr + - ref._sizeBytes / 2)); + checkPointer(ref, i, 1, (ptr + ref._sizeBytes / 2)); if (ref._sizeBytes > 1) { - checkPointer(ref, i, 2, (ptr + - ref._sizeBytes - 1)); + checkPointer(ref, i, 2, (ptr + ref._sizeBytes - 1)); } if (ref._attrib.memoryType == hipMemoryTypeDevice) { - hipFree(ref._pointer); + HIP_CHECK(hipFree(ref._pointer)); } else { - hipHostFree(ref._pointer); + HIP_CHECK(hipHostFree(ref._pointer)); } } } @@ -231,15 +217,13 @@ TEST_CASE("Unit_hipPointerGetAttributes_Basic") { hipError_t e; HIP_CHECK(hipMalloc(&A_d, Nbytes)); - HIP_CHECK(hipHostMalloc(reinterpret_cast(&A_Pinned_h), Nbytes, - hipHostMallocDefault)); + HIP_CHECK(hipHostMalloc(reinterpret_cast(&A_Pinned_h), Nbytes, hipHostMallocDefault)); A_OSAlloc_h = reinterpret_cast(malloc(Nbytes)); size_t free, total; HIP_CHECK(hipMemGetInfo(&free, &total)); printf("hipMemGetInfo: free=%zu (%4.2f) Nbytes=%lu total=%zu (%4.2f)\n", free, - (free / 1024.0 / 1024.0), Nbytes, total, - (total / 1024.0 / 1024.0)); + (free / 1024.0 / 1024.0), Nbytes, total, (total / 1024.0 / 1024.0)); REQUIRE(free + Nbytes <= total); @@ -253,23 +237,20 @@ TEST_CASE("Unit_hipPointerGetAttributes_Basic") { // Check pointer arithmetic cases: resetAttribs(&attribs2); HIP_CHECK(hipPointerGetAttributes(&attribs2, A_d + 100)); - char *ptr = reinterpret_cast(attribs.devicePointer); - REQUIRE(ptr + 100 == - reinterpret_cast(attribs2.devicePointer)); + char* ptr = reinterpret_cast(attribs.devicePointer); + REQUIRE(ptr + 100 == reinterpret_cast(attribs2.devicePointer)); // Corner case at end of array: resetAttribs(&attribs2); HIP_CHECK(hipPointerGetAttributes(&attribs2, A_d + Nbytes - 1)); - REQUIRE((ptr + Nbytes - 1) == - reinterpret_cast(attribs2.devicePointer)); + REQUIRE((ptr + Nbytes - 1) == reinterpret_cast(attribs2.devicePointer)); // Pointer just beyond array must be invalid or at least a different pointer resetAttribs(&attribs2); e = hipPointerGetAttributes(&attribs2, A_d + Nbytes + 1); if (e != hipErrorInvalidValue) { // We might have strayed into another pointer area. - REQUIRE(reinterpret_cast(ptr) != - reinterpret_cast(attribs2.devicePointer)); + REQUIRE(reinterpret_cast(ptr) != reinterpret_cast(attribs2.devicePointer)); } @@ -278,7 +259,7 @@ TEST_CASE("Unit_hipPointerGetAttributes_Basic") { if (e != hipErrorInvalidValue) { REQUIRE(attribs.devicePointer != attribs2.devicePointer); } - hipFree(A_d); + HIP_CHECK(hipFree(A_d)); e = hipPointerGetAttributes(&attribs, A_d); REQUIRE(e == hipErrorInvalidValue); @@ -288,12 +269,11 @@ TEST_CASE("Unit_hipPointerGetAttributes_Basic") { resetAttribs(&attribs2); HIP_CHECK(hipPointerGetAttributes(&attribs2, A_Pinned_h + Nbytes / 2)); - char *ptr1 = reinterpret_cast(attribs.hostPointer); - REQUIRE((ptr1 + Nbytes / 2) - == reinterpret_cast(attribs2.hostPointer)); + char* ptr1 = reinterpret_cast(attribs.hostPointer); + REQUIRE((ptr1 + Nbytes / 2) == reinterpret_cast(attribs2.hostPointer)); - hipHostFree(A_Pinned_h); + HIP_CHECK(hipHostFree(A_Pinned_h)); e = hipPointerGetAttributes(&attribs, A_Pinned_h); REQUIRE(e == hipErrorInvalidValue); @@ -317,33 +297,37 @@ TEST_CASE("Unit_hipPointerGetAttributes_TinyClusterAlloc") { // Multi-threaded test with many simul allocs. // IN : serialize will force the test to run in serial fashion. +#if 0 // FIXME_jatinx These need to be ported to HIP_CHECK_THREAD. Disabling it for now TEST_CASE("Unit_hipPointerGetAttributes_MultiThread") { - srand(0x300); - auto serialize = 1; - printf("\n=============================================\n"); - printf("MultiThreaded_1\n"); - if (serialize) printf("[SERIALIZE]\n"); - printf("===============================================\n"); - std::thread t1(clusterAllocs, 1000, 101, 1000); - if (serialize) t1.join(); + srand(0x300); + auto serialize = 1; + printf("\n=============================================\n"); + printf("MultiThreaded_1\n"); + if (serialize) printf("[SERIALIZE]\n"); + printf("===============================================\n"); + std::thread t1(clusterAllocs, 1000, 101, 1000); + if (serialize) t1.join(); - std::thread t2(clusterAllocs, 1000, 11, 100); - if (serialize) t2.join(); + std::thread t2(clusterAllocs, 1000, 11, 100); + if (serialize) t2.join(); - std::thread t3(clusterAllocs, 1000, 5, 10); - if (serialize) t3.join(); + std::thread t3(clusterAllocs, 1000, 5, 10); + if (serialize) t3.join(); - std::thread t4(clusterAllocs, 1000, 1, 4); - if (serialize) t4.join(); + std::thread t4(clusterAllocs, 1000, 1, 4); + if (serialize) t4.join(); } +#endif TEST_CASE("Unit_hipPointerGetAttributes_Negative") { +#if HT_AMD // Nvidia crashed in hipPointerGetAttributes on nullptr SECTION("Invalid Attributes Pointer") { int* dPtr{nullptr}; HIP_CHECK(hipMalloc(&dPtr, sizeof(int))); HIP_CHECK_ERROR(hipPointerGetAttributes(nullptr, dPtr), hipErrorInvalidValue); HIP_CHECK(hipFree(dPtr)); } +#endif SECTION("Invalid Device Pointer") { hipPointerAttribute_t attributes{}; diff --git a/tests/catch/unit/streamperthread/hipStreamPerThread_Basic.cc b/tests/catch/unit/streamperthread/hipStreamPerThread_Basic.cc index 3780631141..c21c7e25e2 100644 --- a/tests/catch/unit/streamperthread/hipStreamPerThread_Basic.cc +++ b/tests/catch/unit/streamperthread/hipStreamPerThread_Basic.cc @@ -78,8 +78,8 @@ TEST_CASE("Unit_hipStreamPerThread_StreamSynchronize") { constexpr unsigned int MAX_THREAD_CNT = 10; std::vector threads(MAX_THREAD_CNT); - for (auto &th : threads) { - th = std::thread([](){HIP_CHECK(hipStreamSynchronize(hipStreamPerThread));}); + for (auto& th : threads) { + th = std::thread([]() { HIP_CHECK(hipStreamSynchronize(hipStreamPerThread)); }); } for (auto& th : threads) { From 234eb0ccb914e864db6cfc7091679ec322eaed6a Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Thu, 4 Aug 2022 12:24:17 +0530 Subject: [PATCH 8/9] Win: Disable failing hipFree tests --- tests/catch/hipTestMain/config/config_amd_windows.json | 6 +++++- 1 file changed, 5 insertions(+), 1 deletion(-) diff --git a/tests/catch/hipTestMain/config/config_amd_windows.json b/tests/catch/hipTestMain/config/config_amd_windows.json index c0ddfbc0a4..2e5db23232 100644 --- a/tests/catch/hipTestMain/config/config_amd_windows.json +++ b/tests/catch/hipTestMain/config/config_amd_windows.json @@ -54,6 +54,10 @@ "Unit_hipMemGetInfo_Malloc3DArray - char", "Unit_hipMemGetInfo_Malloc3DArray - int", "Unit_hipMemGetInfo_Malloc3DArray - int4", - "Unit_hipMemGetInfo_ParaSmall" + "Unit_hipMemGetInfo_ParaSmall", + "Unit_hipMemGetInfo_ParaMultiSmall", + "Unit_hipFreeMultiTDev", + "Unit_hipFreeMultiTHost", + "Unit_hipFreeMultiTArray" ] } From cfaa115959b9fcf7b120d3bb36fefe227ddc8860 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Thu, 4 Aug 2022 13:23:39 +0530 Subject: [PATCH 9/9] Win: failing sub-tests need to be disabled explicitly --- .../hipTestMain/config/config_amd_windows.json | 15 ++++++++++++--- 1 file changed, 12 insertions(+), 3 deletions(-) diff --git a/tests/catch/hipTestMain/config/config_amd_windows.json b/tests/catch/hipTestMain/config/config_amd_windows.json index 2e5db23232..d96dff5c12 100644 --- a/tests/catch/hipTestMain/config/config_amd_windows.json +++ b/tests/catch/hipTestMain/config/config_amd_windows.json @@ -56,8 +56,17 @@ "Unit_hipMemGetInfo_Malloc3DArray - int4", "Unit_hipMemGetInfo_ParaSmall", "Unit_hipMemGetInfo_ParaMultiSmall", - "Unit_hipFreeMultiTDev", - "Unit_hipFreeMultiTHost", - "Unit_hipFreeMultiTArray" + "Unit_hipFreeMultiTDev - char", + "Unit_hipFreeMultiTDev - int", + "Unit_hipFreeMultiTDev - float2", + "Unit_hipFreeMultiTDev - float4", + "Unit_hipFreeMultiTHost - char", + "Unit_hipFreeMultiTHost - int", + "Unit_hipFreeMultiTHost - float2", + "Unit_hipFreeMultiTHost - float4", + "Unit_hipFreeMultiTArray - char", + "Unit_hipFreeMultiTArray - int", + "Unit_hipFreeMultiTArray - float2", + "Unit_hipFreeMultiTArray - float4" ] }