diff --git a/catch/include/hip_test_defgroups.hh b/catch/include/hip_test_defgroups.hh index d16c5a0d4a..7ade70efd7 100644 --- a/catch/include/hip_test_defgroups.hh +++ b/catch/include/hip_test_defgroups.hh @@ -115,7 +115,14 @@ THE SOFTWARE. * @} */ - /** +/** + * @defgroup PerformanceTest Performance tests + * @{ + * This section describes performance tests for the target API groups and use-cases. + * @} + */ + +/** * @defgroup ShflTest warp shuffle function Management * @{ * This section describes the warp shuffle types & functions of HIP runtime API. diff --git a/catch/include/performance_common.hh b/catch/include/performance_common.hh index 21e87ceb2c..af6ba7609f 100644 --- a/catch/include/performance_common.hh +++ b/catch/include/performance_common.hh @@ -34,6 +34,7 @@ THE SOFTWARE. #include #pragma clang diagnostic ignored "-Wunused-but-set-variable" +#pragma clang diagnostic ignored "-Wunused-parameter" #pragma clang diagnostic ignored "-Wunused-function" #if defined(_WIN32) diff --git a/catch/performance/CMakeLists.txt b/catch/performance/CMakeLists.txt index e1e159c6e5..0c6962c596 100644 --- a/catch/performance/CMakeLists.txt +++ b/catch/performance/CMakeLists.txt @@ -18,5 +18,6 @@ # OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN # THE SOFTWARE. +add_subdirectory(stream) add_subdirectory(event) add_subdirectory(example) diff --git a/catch/performance/stream/CMakeLists.txt b/catch/performance/stream/CMakeLists.txt new file mode 100644 index 0000000000..f9504d63da --- /dev/null +++ b/catch/performance/stream/CMakeLists.txt @@ -0,0 +1,63 @@ +# 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. + +if(HIP_PLATFORM MATCHES "amd") +set(TEST_SRC + hipStreamWaitEvent.cc + hipStreamGetFlags.cc + hipStreamGetPriority.cc + hipExtStreamCreateWithCUMask.cc + hipExtStreamGetCUMask.cc + hipStreamAddCallback.cc + hipStreamWaitValue.cc + hipStreamWriteValue.cc + hipMallocAsync.cc + hipFreeAsync.cc + hipMemPoolCreate.cc + hipMemPoolDestroy.cc + hipMemPoolTrimTo.cc + hipMemPoolSetAttribute.cc + hipMemPoolGetAttribute.cc + hipMemPoolSetAccess.cc + hipMallocFromPoolAsync.cc + hipMemPoolExportToShareableHandle.cc + hipMemPoolImportFromShareableHandle.cc + hipMemPoolExportPointer.cc + hipMemPoolImportPointer.cc + hipStreamBasic.cc +) +else() +set(TEST_SRC + hipStreamWaitEvent.cc + hipStreamGetFlags.cc + hipStreamGetPriority.cc + hipStreamAddCallback.cc + hipStreamWaitValue.cc + hipStreamWriteValue.cc + hipMallocAsync.cc + hipFreeAsync.cc + hipStreamBasic.cc +) +endif() + +hip_add_exe_to_target(NAME StreamPerformance + TEST_SRC ${TEST_SRC} + TEST_TARGET_NAME build_tests + COMPILE_OPTIONS -std=c++17) diff --git a/catch/performance/stream/hipExtStreamCreateWithCUMask.cc b/catch/performance/stream/hipExtStreamCreateWithCUMask.cc new file mode 100644 index 0000000000..f09ae8e976 --- /dev/null +++ b/catch/performance/stream/hipExtStreamCreateWithCUMask.cc @@ -0,0 +1,65 @@ +/* +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 + +/** + * @addtogroup stream stream + * @{ + * @ingroup PerformanceTest + * Contains performance tests for all stream management HIP APIs. + */ + +class ExtStreamCreateWithCUMaskBenchmark : public Benchmark { + public: + void operator()() { + hipDeviceProp_t props; + HIP_CHECK(hipGetDeviceProperties(&props, 0)); + std::vector cu_mask(props.multiProcessorCount, 0); + hipStream_t stream{}; + + TIMED_SECTION(kTimerTypeCpu) { + HIP_CHECK(hipExtStreamCreateWithCUMask(&stream, cu_mask.size(), cu_mask.data())); + } + + HIP_CHECK(hipStreamDestroy(stream)); + } +}; + +static void RunBenchmark() { + ExtStreamCreateWithCUMaskBenchmark benchmark; + benchmark.Run(); +} + +/** + * Test Description + * ------------------------ + * - Executes `hipExtStreamCreateWithCUMask`. + * Test source + * ------------------------ + * - performance/stream/hipExtStreamCreateWithCUMask.cc + * Test requirements + * ------------------------ + * - Platform specific (AMD) + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Performance_hipExtStreamCreateWithCUMask") { + RunBenchmark(); +} diff --git a/catch/performance/stream/hipExtStreamGetCUMask.cc b/catch/performance/stream/hipExtStreamGetCUMask.cc new file mode 100644 index 0000000000..1dd06aaed6 --- /dev/null +++ b/catch/performance/stream/hipExtStreamGetCUMask.cc @@ -0,0 +1,67 @@ +/* +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 + +/** + * @addtogroup stream stream + * @{ + * @ingroup PerformanceTest + */ + +class ExtStreamGetCUMaskBenchmark : public Benchmark { + public: + void operator()() { + hipDeviceProp_t props; + HIP_CHECK(hipGetDeviceProperties(&props, 0)); + std::vector cu_mask(props.multiProcessorCount, 0); + hipStream_t stream{}; + HIP_CHECK(hipExtStreamCreateWithCUMask(&stream, cu_mask.size(), cu_mask.data())); + std::vector new_cu_mask(cu_mask.size(), 0); + + TIMED_SECTION(kTimerTypeCpu) { + HIP_CHECK(hipExtStreamGetCUMask(stream, new_cu_mask.size(), new_cu_mask.data())); + } + + HIP_CHECK(hipStreamDestroy(stream)); + } +}; + +static void RunBenchmark() { + ExtStreamGetCUMaskBenchmark benchmark; + benchmark.Run(); +} + +/** + * Test Description + * ------------------------ + * - Executes `hipExtStreamGetCUMask`. + * - Creates basic mask and gets it into the new one. + * Test source + * ------------------------ + * - performance/stream/hipExtStreamGetCUMask.cc + * Test requirements + * ------------------------ + * - Platform specific (AMD) + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Performance_hipExtStreamGetCUMask") { + RunBenchmark(); +} diff --git a/catch/performance/stream/hipFreeAsync.cc b/catch/performance/stream/hipFreeAsync.cc new file mode 100644 index 0000000000..0e21ac4b2a --- /dev/null +++ b/catch/performance/stream/hipFreeAsync.cc @@ -0,0 +1,69 @@ +/* +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 + +/** + * @addtogroup stream stream + * @{ + * @ingroup PerformanceTest + */ + +class FreeAsyncBenchmark : public Benchmark { + public: + void operator()(const size_t array_size) { + const StreamGuard stream_guard{Streams::created}; + const hipStream_t stream = stream_guard.stream(); + float* dev_ptr{nullptr}; + HIP_CHECK(hipMallocAsync(reinterpret_cast(&dev_ptr), array_size * sizeof(float), stream)); + + TIMED_SECTION_STREAM(kTimerTypeEvent, stream) { + HIP_CHECK(hipFreeAsync(dev_ptr, stream)); + } + + HIP_CHECK(hipStreamSynchronize(stream)); + } +}; + +static void RunBenchmark(const size_t array_size) { + FreeAsyncBenchmark benchmark; + benchmark.AddSectionName(std::to_string(array_size)); + benchmark.Run(array_size); +} + +/** + * Test Description + * ------------------------ + * - Executes `hipFreeAsync` with created stream: + * -# Allocation size: + * - 4 KB + * - 4 MB + * - 16 MB + * Test source + * ------------------------ + * - performance/stream/hipFreeAsync.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Performance_hipFreeAsync") { + size_t array_size = GENERATE(4_KB, 4_MB, 16_MB); + RunBenchmark(array_size); +} diff --git a/catch/performance/stream/hipMallocAsync.cc b/catch/performance/stream/hipMallocAsync.cc new file mode 100644 index 0000000000..8bae9e0bfe --- /dev/null +++ b/catch/performance/stream/hipMallocAsync.cc @@ -0,0 +1,68 @@ +/* +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 + +/** + * @addtogroup stream stream + * @{ + * @ingroup PerformanceTest + */ + +class MallocAsyncBenchmark : public Benchmark { + public: + void operator()(const size_t array_size) { + const StreamGuard stream_guard{Streams::created}; + const hipStream_t stream = stream_guard.stream(); + float* dev_ptr{nullptr}; + + TIMED_SECTION_STREAM(kTimerTypeEvent, stream) { + HIP_CHECK(hipMallocAsync(reinterpret_cast(&dev_ptr), array_size * sizeof(float), stream)); + } + HIP_CHECK(hipStreamSynchronize(stream)); + HIP_CHECK(hipFree(dev_ptr)); + } +}; + +static void RunBenchmark(const size_t array_size) { + MallocAsyncBenchmark benchmark; + benchmark.AddSectionName(std::to_string(array_size)); + benchmark.Run(array_size); +} + +/** + * Test Description + * ------------------------ + * - Executes `hipMallocAsync` with created stream: + * -# Allocation size: + * - 4 KB + * - 4 MB + * - 16 MB + * Test source + * ------------------------ + * - performance/stream/hipMallocAsync.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Performance_hipMallocAsync") { + size_t array_size = GENERATE(4_KB, 4_MB, 16_MB); + RunBenchmark(array_size); +} diff --git a/catch/performance/stream/hipMallocFromPoolAsync.cc b/catch/performance/stream/hipMallocFromPoolAsync.cc new file mode 100644 index 0000000000..a0de4bf908 --- /dev/null +++ b/catch/performance/stream/hipMallocFromPoolAsync.cc @@ -0,0 +1,82 @@ +/* +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 "mem_pools_performance_common.hh" + +/** + * @addtogroup stream stream + * @{ + * @ingroup PerformanceTest + */ + +class MallocFromPoolAsyncBenchmark : public Benchmark { + public: + void operator()(const size_t array_size) { + const StreamGuard stream_guard{Streams::created}; + const hipStream_t stream = stream_guard.stream(); + + hipMemPool_t mem_pool{nullptr}; + hipMemPoolProps pool_props = CreateMemPoolProps(0, hipMemHandleTypeNone); + HIP_CHECK(hipMemPoolCreate(&mem_pool, &pool_props)); + + float* array_ptr{nullptr}; + + TIMED_SECTION_STREAM(kTimerTypeEvent, stream) { + HIP_CHECK(hipMallocFromPoolAsync(&array_ptr, array_size * sizeof(float), mem_pool, stream)); + } + + REQUIRE(array_ptr != nullptr); + + HIP_CHECK(hipFreeAsync(array_ptr, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + HIP_CHECK(hipMemPoolDestroy(mem_pool)); + } +}; + +static void RunBenchmark(const size_t array_size) { + MallocFromPoolAsyncBenchmark benchmark; + benchmark.AddSectionName(std::to_string(array_size)); + benchmark.Run(array_size); +} + +/** + * Test Description + * ------------------------ + * - Executes `hipMallocFromPoolAsync`: + * -# Allocation size: + * - 4 KB + * - 4 MB + * - 16 MB + * Test source + * ------------------------ + * - performance/stream/hipMallocFromPoolAsync.cc + * Test requirements + * ------------------------ + * - Device supports memory pools + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Performance_hipMallocFromPoolAsync") { + if (!AreMemPoolsSupported(0)) { + HipTest::HIP_SKIP_TEST("GPU 0 doesn't support hipDeviceAttributeMemoryPoolsSupported " + "attribute. Hence skipping the testing with Pass result.\n"); + return; + } + size_t array_size = GENERATE(4_KB, 4_MB, 16_MB); + RunBenchmark(array_size); +} diff --git a/catch/performance/stream/hipMemPoolCreate.cc b/catch/performance/stream/hipMemPoolCreate.cc new file mode 100644 index 0000000000..3e616b1991 --- /dev/null +++ b/catch/performance/stream/hipMemPoolCreate.cc @@ -0,0 +1,71 @@ +/* +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 "mem_pools_performance_common.hh" + +/** + * @addtogroup stream stream + * @{ + * @ingroup PerformanceTest + */ + +class MemPoolCreateBenchmark : public Benchmark { + public: + void operator()() { + hipMemPool_t mem_pool{nullptr}; + hipMemPoolProps pool_props = CreateMemPoolProps(0, hipMemHandleTypeNone); + + TIMED_SECTION(kTimerTypeCpu) { + HIP_CHECK(hipMemPoolCreate(&mem_pool, &pool_props)); + } + + REQUIRE(mem_pool != nullptr); + HIP_CHECK(hipMemPoolDestroy(mem_pool)); + } +}; + +static void RunBenchmark() { + MemPoolCreateBenchmark benchmark; + benchmark.Run(); +} + +/** + * @warning **MemPool APIs are not fully implemented within current version + * or HIP and therefore they cannot be appropriately executed on AMD and NVIDIA platforms. + * Therefore, all tests related to MemPool APIs are implemented without formal + * verification and will be verified once HIP fully supports MemPool APIs.** + * Test Description + * ------------------------ + * - Executes `hipMemPoolCreate`. + * Test source + * ------------------------ + * - performance/stream/hipMemPoolCreate.cc + * Test requirements + * ------------------------ + * - Device supports memory pools + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Performance_hipMemPoolCreate") { + if (!AreMemPoolsSupported(0)) { + HipTest::HIP_SKIP_TEST("GPU 0 doesn't support hipDeviceAttributeMemoryPoolsSupported " + "attribute. Hence skipping the testing with Pass result.\n"); + return; + } + RunBenchmark(); +} diff --git a/catch/performance/stream/hipMemPoolDestroy.cc b/catch/performance/stream/hipMemPoolDestroy.cc new file mode 100644 index 0000000000..0276ccc962 --- /dev/null +++ b/catch/performance/stream/hipMemPoolDestroy.cc @@ -0,0 +1,70 @@ +/* +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 "mem_pools_performance_common.hh" + +/** + * @addtogroup stream stream + * @{ + * @ingroup PerformanceTest + */ + +class MemPoolDestroyBenchmark : public Benchmark { + public: + void operator()() { + hipMemPool_t mem_pool{nullptr}; + hipMemPoolProps pool_props = CreateMemPoolProps(0, hipMemHandleTypeNone); + HIP_CHECK(hipMemPoolCreate(&mem_pool, &pool_props)); + + TIMED_SECTION(kTimerTypeCpu) { + HIP_CHECK(hipMemPoolDestroy(mem_pool)); + } + } +}; + +static void RunBenchmark() { + MemPoolDestroyBenchmark benchmark; + benchmark.Run(); +} + +/** + * @warning **MemPool APIs are not fully implemented within current version + * or HIP and therefore they cannot be appropriately executed on AMD and NVIDIA platforms. + * Therefore, all tests related to MemPool APIs are implemented without formal + * verification and will be verified once HIP fully supports MemPool APIs.** + * Test Description + * ------------------------ + * - Creates new mem pool. + * - Executes `hipMemPoolDestroy`. + * Test source + * ------------------------ + * - performance/stream/hipMemPoolDestroy.cc + * Test requirements + * ------------------------ + * - Device supports memory pools + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Performance_hipMemPoolDestroy") { + if (!AreMemPoolsSupported(0)) { + HipTest::HIP_SKIP_TEST("GPU 0 doesn't support hipDeviceAttributeMemoryPoolsSupported " + "attribute. Hence skipping the testing with Pass result.\n"); + return; + } + RunBenchmark(); +} diff --git a/catch/performance/stream/hipMemPoolExportPointer.cc b/catch/performance/stream/hipMemPoolExportPointer.cc new file mode 100644 index 0000000000..10960706d0 --- /dev/null +++ b/catch/performance/stream/hipMemPoolExportPointer.cc @@ -0,0 +1,84 @@ +/* +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 "mem_pools_performance_common.hh" + +/** + * @addtogroup stream stream + * @{ + * @ingroup PerformanceTest + */ + +class MemPoolExportPointerBenchmark : public Benchmark { + public: + void operator()(const size_t array_size) { + float* device_ptr{nullptr}; + hipMemPool_t mem_pool{nullptr}; + hipMemPoolPtrExportData exp_data; + + hipMemPoolProps props = CreateMemPoolProps(0, kHandleType); + HIP_CHECK(hipMemPoolCreate(&mem_pool, &props)); + HIP_CHECK(hipMallocFromPoolAsync(&device_ptr, array_size * sizeof(float), mem_pool, nullptr)); + HIP_CHECK(hipStreamSynchronize(nullptr)); + + TIMED_SECTION(kTimerTypeCpu) { + HIP_CHECK(hipMemPoolExportPointer(&exp_data, device_ptr)); + } + + HIP_CHECK(hipFreeAsync(device_ptr, nullptr)); + HIP_CHECK(hipMemPoolDestroy(mem_pool)); + } +}; + +static void RunBenchmark(const size_t array_size) { + MemPoolExportPointerBenchmark benchmark; + benchmark.AddSectionName(std::to_string(array_size)); + benchmark.Run(array_size); +} + +/** + * @warning **MemPool APIs are not fully implemented within current version + * or HIP and therefore they cannot be appropriately executed on AMD and NVIDIA platforms. + * Therefore, all tests related to MemPool APIs are implemented without formal + * verification and will be verified once HIP fully supports MemPool APIs.** + * Test Description + * ------------------------ + * - Executes `hipMemPoolExportPointer`: + * -# Allocation size: + * - 4 KB + * - 4 MB + * - 16 MB + * - Uses the same process for import and export operations. + * Test source + * ------------------------ + * - performance/stream/hipMemPoolExportPointer.cc + * Test requirements + * ------------------------ + * - Device supports memory pools + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Performance_hipMemPoolExportPointer") { + if (!AreMemPoolsSupported(0)) { + HipTest::HIP_SKIP_TEST("GPU 0 doesn't support hipDeviceAttributeMemoryPoolsSupported " + "attribute. Hence skipping the testing with Pass result.\n"); + return; + } + size_t array_size = GENERATE(4_KB, 4_MB, 16_MB); + RunBenchmark(array_size); +} diff --git a/catch/performance/stream/hipMemPoolExportToShareableHandle.cc b/catch/performance/stream/hipMemPoolExportToShareableHandle.cc new file mode 100644 index 0000000000..9e93751403 --- /dev/null +++ b/catch/performance/stream/hipMemPoolExportToShareableHandle.cc @@ -0,0 +1,74 @@ +/* +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 "mem_pools_performance_common.hh" + +/** + * @addtogroup stream stream + * @{ + * @ingroup PerformanceTest + */ + +class MemPoolExportToShareableHandleBenchmark : public Benchmark { + public: + void operator()() { + hipMemPool_t mem_pool{nullptr}; + int share_handle; + + hipMemPoolProps props = CreateMemPoolProps(0, kHandleType); + HIP_CHECK(hipMemPoolCreate(&mem_pool, &props)); + + TIMED_SECTION(kTimerTypeCpu) { + HIP_CHECK(hipMemPoolExportToShareableHandle(&share_handle, mem_pool, kHandleType, 0)); + } + + HIP_CHECK(hipMemPoolDestroy(mem_pool)); + } +}; + +static void RunBenchmark() { + MemPoolExportToShareableHandleBenchmark benchmark; + benchmark.Run(); +} + +/** + * @warning **MemPool APIs are not fully implemented within current version + * or HIP and therefore they cannot be appropriately executed on AMD and NVIDIA platforms. + * Therefore, all tests related to MemPool APIs are implemented without formal + * verification and will be verified once HIP fully supports MemPool APIs.** + * Test Description + * ------------------------ + * - Executes `hipMemPoolExportToShareableHandle`. + * - Uses the same process for import and export operations. + * Test source + * ------------------------ + * - performance/stream/hipMemPoolExportToShareableHandle.cc + * Test requirements + * ------------------------ + * - Device supports memory pools + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Performance_hipMemPoolExportToShareableHandle") { + if (!AreMemPoolsSupported(0)) { + HipTest::HIP_SKIP_TEST("GPU 0 doesn't support hipDeviceAttributeMemoryPoolsSupported " + "attribute. Hence skipping the testing with Pass result.\n"); + return; + } + RunBenchmark(); +} diff --git a/catch/performance/stream/hipMemPoolGetAccess.cc b/catch/performance/stream/hipMemPoolGetAccess.cc new file mode 100644 index 0000000000..c30b45acb6 --- /dev/null +++ b/catch/performance/stream/hipMemPoolGetAccess.cc @@ -0,0 +1,76 @@ +/* +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 "mem_pools_performance_common.hh" + +/** + * @addtogroup stream stream + * @{ + * @ingroup PerformanceTest + */ + +class MemPoolGetAccessBenchmark : public Benchmark { + public: + void operator()() { + hipMemPool_t mem_pool{nullptr}; + hipMemPoolProps pool_props = CreateMemPoolProps(0, hipMemHandleTypeNone); + HIP_CHECK(hipMemPoolCreate(&mem_pool, &pool_props)); + + hipMemAccessFlags flags = hipMemAccessFlagsProtNone; + hipMemLocation location = { + hipMemLocationTypeDevice, + 0 + }; + TIMED_SECTION(kTimerTypeCpu) { + HIP_CHECK(hipMemPoolGetAccess(&flags, mem_pool, location)); + } + + HIP_CHECK(hipMemPoolDestroy(mem_pool)); + } +}; + +static void RunBenchmark() { + MemPoolGetAccessBenchmark benchmark; + benchmark.Run(); +} + +/** + * @warning **MemPool APIs are not fully implemented within current version + * or HIP and therefore they cannot be appropriately executed on AMD and NVIDIA platforms. + * Therefore, all tests related to MemPool APIs are implemented without formal + * verification and will be verified once HIP fully supports MemPool APIs.** + * Test Description + * ------------------------ + * - Executes `hipMemPoolGetAccess`. + * Test source + * ------------------------ + * - performance/stream/hipMemPoolGetAccess.cc + * Test requirements + * ------------------------ + * - Device supports memory pools + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Performance_hipMemPoolGetAccess") { + if (!AreMemPoolsSupported(0)) { + HipTest::HIP_SKIP_TEST("GPU 0 doesn't support hipDeviceAttributeMemoryPoolsSupported " + "attribute. Hence skipping the testing with Pass result.\n"); + return; + } + RunBenchmark(); +} diff --git a/catch/performance/stream/hipMemPoolGetAttribute.cc b/catch/performance/stream/hipMemPoolGetAttribute.cc new file mode 100644 index 0000000000..d034fac57c --- /dev/null +++ b/catch/performance/stream/hipMemPoolGetAttribute.cc @@ -0,0 +1,83 @@ +/* +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 "mem_pools_performance_common.hh" + +/** + * @addtogroup stream stream + * @{ + * @ingroup PerformanceTest + */ + +class MemPoolGetAttributeBenchmark : public Benchmark { + public: + void operator()(const hipMemPoolAttr attribute) { + hipMemPool_t mem_pool{nullptr}; + hipMemPoolProps pool_props = CreateMemPoolProps(0, hipMemHandleTypeNone); + HIP_CHECK(hipMemPoolCreate(&mem_pool, &pool_props)); + + uint64_t value{0}; + + TIMED_SECTION(kTimerTypeCpu) { + HIP_CHECK(hipMemPoolGetAttribute(mem_pool, attribute, &value)); + } + + HIP_CHECK(hipMemPoolDestroy(mem_pool)); + } +}; + +static void RunBenchmark(const hipMemPoolAttr attribute) { + MemPoolGetAttributeBenchmark benchmark; + benchmark.AddSectionName(GetMemPoolAttrSectionName(attribute)); + benchmark.Run(attribute); +} + +/** + * @warning **MemPool APIs are not fully implemented within current version + * or HIP and therefore they cannot be appropriately executed on AMD and NVIDIA platforms. + * Therefore, all tests related to MemPool APIs are implemented without formal + * verification and will be verified once HIP fully supports MemPool APIs.** + * Test Description + * ------------------------ + * - Executes `hipMemPoolGetAttribute`: + * -# Supported attributes: + * - `hipMemPoolAttrReleaseThreshold` + * - `hipMemPoolReuseFollowEventDependencies` + * - `hipMemPoolReuseAllowOpportunistic` + * - `hipMemPoolReuseAllowInternalDependencies` + * Test source + * ------------------------ + * - performance/stream/hipMemPoolGetAttribute.cc + * Test requirements + * ------------------------ + * - Device supports memory pools + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Performance_hipMemPoolGetAttribute") { + if (!AreMemPoolsSupported(0)) { + HipTest::HIP_SKIP_TEST("GPU 0 doesn't support hipDeviceAttributeMemoryPoolsSupported " + "attribute. Hence skipping the testing with Pass result.\n"); + return; + } + hipMemPoolAttr attribute = GENERATE(hipMemPoolAttrReleaseThreshold, + hipMemPoolReuseFollowEventDependencies, + hipMemPoolReuseAllowOpportunistic, + hipMemPoolReuseAllowInternalDependencies); + RunBenchmark(attribute); +} diff --git a/catch/performance/stream/hipMemPoolImportFromShareableHandle.cc b/catch/performance/stream/hipMemPoolImportFromShareableHandle.cc new file mode 100644 index 0000000000..fb37f6c951 --- /dev/null +++ b/catch/performance/stream/hipMemPoolImportFromShareableHandle.cc @@ -0,0 +1,75 @@ +/* +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 "mem_pools_performance_common.hh" + +/** + * @addtogroup stream stream + * @{ + * @ingroup PerformanceTest + */ + +class MemPoolImportFromShareableHandleBenchmark : public Benchmark { + public: + void operator()() { + hipMemPool_t mem_pool{nullptr}; + int share_handle; + + hipMemPoolProps props = CreateMemPoolProps(0, kHandleType); + HIP_CHECK(hipMemPoolCreate(&mem_pool, &props)); + HIP_CHECK(hipMemPoolExportToShareableHandle(&share_handle, mem_pool, kHandleType, 0)); + + TIMED_SECTION(kTimerTypeCpu) { + HIP_CHECK(hipMemPoolImportFromShareableHandle(&mem_pool, &share_handle, kHandleType, 0)); + } + + HIP_CHECK(hipMemPoolDestroy(mem_pool)); + } +}; + +static void RunBenchmark() { + MemPoolImportFromShareableHandleBenchmark benchmark; + benchmark.Run(); +} + +/** + * @warning **MemPool APIs are not fully implemented within current version + * or HIP and therefore they cannot be appropriately executed on AMD and NVIDIA platforms. + * Therefore, all tests related to MemPool APIs are implemented without formal + * verification and will be verified once HIP fully supports MemPool APIs.** + * Test Description + * ------------------------ + * - Executes `hipMemPoolImportFromShareableHandle`. + * - Uses the same process for import and export operations. + * Test source + * ------------------------ + * - performance/stream/hipMemPoolImportFromShareableHandle.cc + * Test requirements + * ------------------------ + * - Device supports memory pools + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Performance_hipMemPoolImportFromShareableHandle") { + if (!AreMemPoolsSupported(0)) { + HipTest::HIP_SKIP_TEST("GPU 0 doesn't support hipDeviceAttributeMemoryPoolsSupported " + "attribute. Hence skipping the testing with Pass result.\n"); + return; + } + RunBenchmark(); +} diff --git a/catch/performance/stream/hipMemPoolImportPointer.cc b/catch/performance/stream/hipMemPoolImportPointer.cc new file mode 100644 index 0000000000..de439e7c93 --- /dev/null +++ b/catch/performance/stream/hipMemPoolImportPointer.cc @@ -0,0 +1,87 @@ +/* +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 "mem_pools_performance_common.hh" + +/** + * @addtogroup stream stream + * @{ + * @ingroup PerformanceTest + */ + +class MemPoolImportPointerBenchmark : public Benchmark { + public: + void operator()(const size_t array_size) { + float* device_ptr{nullptr}; + float* device_ptr_import{nullptr}; + hipMemPool_t mem_pool{nullptr}; + hipMemPoolPtrExportData exp_data; + + hipMemPoolProps props = CreateMemPoolProps(0, kHandleType); + HIP_CHECK(hipMemPoolCreate(&mem_pool, &props)); + HIP_CHECK(hipMallocFromPoolAsync(&device_ptr, array_size * sizeof(float), mem_pool, nullptr)); + HIP_CHECK(hipStreamSynchronize(nullptr)); + HIP_CHECK(hipMemPoolExportPointer(&exp_data, device_ptr)); + + TIMED_SECTION(kTimerTypeCpu) { + HIP_CHECK(hipMemPoolImportPointer(reinterpret_cast(device_ptr_import), mem_pool, &exp_data)); + } + + HIP_CHECK(hipFree(device_ptr)); + HIP_CHECK(hipFree(device_ptr_import)); + HIP_CHECK(hipMemPoolDestroy(mem_pool)); + } +}; + +static void RunBenchmark(const size_t array_size) { + MemPoolImportPointerBenchmark benchmark; + benchmark.AddSectionName(std::to_string(array_size)); + benchmark.Run(array_size); +} + +/** + * @warning **MemPool APIs are not fully implemented within current version + * or HIP and therefore they cannot be appropriately executed on AMD and NVIDIA platforms. + * Therefore, all tests related to MemPool APIs are implemented without formal + * verification and will be verified once HIP fully supports MemPool APIs.** + * Test Description + * ------------------------ + * - Executes `hipMemPoolImportPointer`: + * -# Allocation size: + * - 4 KB + * - 4 MB + * - 16 MB + * - Uses the same process for import and export operations. + * Test source + * ------------------------ + * - performance/stream/hipMemPoolImportPointer.cc + * Test requirements + * ------------------------ + * - Device supports memory pools + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Performance_hipMemPoolImportPointer") { + if (!AreMemPoolsSupported(0)) { + HipTest::HIP_SKIP_TEST("GPU 0 doesn't support hipDeviceAttributeMemoryPoolsSupported " + "attribute. Hence skipping the testing with Pass result.\n"); + return; + } + size_t array_size = GENERATE(4_KB, 4_MB, 16_MB); + RunBenchmark(array_size); +} diff --git a/catch/performance/stream/hipMemPoolSetAccess.cc b/catch/performance/stream/hipMemPoolSetAccess.cc new file mode 100644 index 0000000000..7eadcb4568 --- /dev/null +++ b/catch/performance/stream/hipMemPoolSetAccess.cc @@ -0,0 +1,79 @@ +/* +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 "mem_pools_performance_common.hh" + +/** + * @addtogroup stream stream + * @{ + * @ingroup PerformanceTest + */ + +class MemPoolSetAccessBenchmark : public Benchmark { + public: + void operator()() { + hipMemPool_t mem_pool{nullptr}; + hipMemPoolProps pool_props = CreateMemPoolProps(0, hipMemHandleTypeNone); + HIP_CHECK(hipMemPoolCreate(&mem_pool, &pool_props)); + + hipMemAccessDesc desc_list = { + { + hipMemLocationTypeDevice, + 0 + }, + hipMemAccessFlagsProtReadWrite + }; + + TIMED_SECTION(kTimerTypeCpu) { + HIP_CHECK(hipMemPoolSetAccess(mem_pool, &desc_list, 1)); + } + + HIP_CHECK(hipMemPoolDestroy(mem_pool)); + } +}; + +static void RunBenchmark() { + MemPoolSetAccessBenchmark benchmark; + benchmark.Run(); +} + +/** + * @warning **MemPool APIs are not fully implemented within current version + * or HIP and therefore they cannot be appropriately executed on AMD and NVIDIA platforms. + * Therefore, all tests related to MemPool APIs are implemented without formal + * verification and will be verified once HIP fully supports MemPool APIs.** + * Test Description + * ------------------------ + * - Executes `hipMemPoolSetAccess` with `hipMemAccessFlagsProtReadWrite`. + * Test source + * ------------------------ + * - performance/stream/hipMemPoolSetAccess.cc + * Test requirements + * ------------------------ + * - Device supports memory pools + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Performance_hipMemPoolSetAccess") { + if (!AreMemPoolsSupported(0)) { + HipTest::HIP_SKIP_TEST("GPU 0 doesn't support hipDeviceAttributeMemoryPoolsSupported " + "attribute. Hence skipping the testing with Pass result.\n"); + return; + } + RunBenchmark(); +} diff --git a/catch/performance/stream/hipMemPoolSetAttribute.cc b/catch/performance/stream/hipMemPoolSetAttribute.cc new file mode 100644 index 0000000000..dbdc2ca152 --- /dev/null +++ b/catch/performance/stream/hipMemPoolSetAttribute.cc @@ -0,0 +1,83 @@ +/* +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 "mem_pools_performance_common.hh" + +/** + * @addtogroup stream stream + * @{ + * @ingroup PerformanceTest + */ + +class MemPoolSetAttributeBenchmark : public Benchmark { + public: + void operator()(const hipMemPoolAttr attribute) { + hipMemPool_t mem_pool{nullptr}; + hipMemPoolProps pool_props = CreateMemPoolProps(0, hipMemHandleTypeNone); + HIP_CHECK(hipMemPoolCreate(&mem_pool, &pool_props)); + + int value{0}; + + TIMED_SECTION(kTimerTypeCpu) { + HIP_CHECK(hipMemPoolSetAttribute(mem_pool, attribute, &value)); + } + + HIP_CHECK(hipMemPoolDestroy(mem_pool)); + } +}; + +static void RunBenchmark(const hipMemPoolAttr attribute) { + MemPoolSetAttributeBenchmark benchmark; + benchmark.AddSectionName(GetMemPoolAttrSectionName(attribute)); + benchmark.Run(attribute); +} + +/** + * @warning **MemPool APIs are not fully implemented within current version + * or HIP and therefore they cannot be appropriately executed on AMD and NVIDIA platforms. + * Therefore, all tests related to MemPool APIs are implemented without formal + * verification and will be verified once HIP fully supports MemPool APIs.** + * Test Description + * ------------------------ + * - Executes `hipMemPoolSetAttribute`: + * -# Supported attributes: + * - `hipMemPoolAttrReleaseThreshold` + * - `hipMemPoolReuseFollowEventDependencies` + * - `hipMemPoolReuseAllowOpportunistic` + * - `hipMemPoolReuseAllowInternalDependencies` + * Test source + * ------------------------ + * - performance/stream/hipMemPoolSetAttribute.cc + * Test requirements + * ------------------------ + * - Device supports memory pools + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Performance_hipMemPoolSetAttribute") { + if (!AreMemPoolsSupported(0)) { + HipTest::HIP_SKIP_TEST("GPU 0 doesn't support hipDeviceAttributeMemoryPoolsSupported " + "attribute. Hence skipping the testing with Pass result.\n"); + return; + } + hipMemPoolAttr attribute = GENERATE(hipMemPoolAttrReleaseThreshold, + hipMemPoolReuseFollowEventDependencies, + hipMemPoolReuseAllowOpportunistic, + hipMemPoolReuseAllowInternalDependencies); + RunBenchmark(attribute); +} diff --git a/catch/performance/stream/hipMemPoolTrimTo.cc b/catch/performance/stream/hipMemPoolTrimTo.cc new file mode 100644 index 0000000000..e3fc5b6efb --- /dev/null +++ b/catch/performance/stream/hipMemPoolTrimTo.cc @@ -0,0 +1,77 @@ +/* +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 "mem_pools_performance_common.hh" + +/** + * @addtogroup stream stream + * @{ + * @ingroup PerformanceTest + */ + +class MemPoolTrimToBenchmark : public Benchmark { + public: + void operator()(const size_t min_bytes_to_hold) { + hipMemPool_t mem_pool{nullptr}; + hipMemPoolProps pool_props = CreateMemPoolProps(0, hipMemHandleTypeNone); + HIP_CHECK(hipMemPoolCreate(&mem_pool, &pool_props)); + + TIMED_SECTION(kTimerTypeCpu) { + HIP_CHECK(hipMemPoolTrimTo(mem_pool, min_bytes_to_hold)); + } + + HIP_CHECK(hipMemPoolDestroy(mem_pool)); + } +}; + +static void RunBenchmark(const size_t min_bytes_to_hold) { + MemPoolTrimToBenchmark benchmark; + benchmark.AddSectionName(std::to_string(min_bytes_to_hold)); + benchmark.Run(min_bytes_to_hold); +} + +/** + * @warning **MemPool APIs are not fully implemented within current version + * or HIP and therefore they cannot be appropriately executed on AMD and NVIDIA platforms. + * Therefore, all tests related to MemPool APIs are implemented without formal + * verification and will be verified once HIP fully supports MemPool APIs.** + * Test Description + * ------------------------ + * - Executes `hipMemPoolTrimTo`: + * -# Minimum bytes to hold: + * - 4 KB + * - 4 MB + * - 16 MB + * Test source + * ------------------------ + * - performance/stream/hipMemPoolTrimTo.cc + * Test requirements + * ------------------------ + * - Device supports memory pools + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Performance_hipMemPoolTrimTo") { + if (!AreMemPoolsSupported(0)) { + HipTest::HIP_SKIP_TEST("GPU 0 doesn't support hipDeviceAttributeMemoryPoolsSupported " + "attribute. Hence skipping the testing with Pass result.\n"); + return; + } + size_t min_bytes_to_hold = GENERATE(4_KB, 4_MB, 16_MB); + RunBenchmark(min_bytes_to_hold); +} diff --git a/catch/performance/stream/hipStreamAddCallback.cc b/catch/performance/stream/hipStreamAddCallback.cc new file mode 100644 index 0000000000..771d0ee43e --- /dev/null +++ b/catch/performance/stream/hipStreamAddCallback.cc @@ -0,0 +1,61 @@ +/* +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 + +/** + * @addtogroup stream stream + * @{ + * @ingroup PerformanceTest + */ + +void Callback(hipStream_t stream, hipError_t status, void* user_data) {} + +class StreamAddCallbackBenchmark : public Benchmark { + public: + void operator()() { + const StreamGuard stream_guard{Streams::created}; + const hipStream_t stream = stream_guard.stream(); + + TIMED_SECTION(kTimerTypeCpu) { + HIP_CHECK(hipStreamAddCallback(stream, Callback, nullptr, 0)); + } + } +}; + +static void RunBenchmark() { + StreamAddCallbackBenchmark benchmark; + benchmark.Run(); +} + +/** + * Test Description + * ------------------------ + * - Executes `hipStreamAddCallback` on the created stream. + * Test source + * ------------------------ + * - performance/stream/hipStreamAddCallback.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Performance_hipStreamAddCallback") { + RunBenchmark(); +} diff --git a/catch/performance/stream/hipStreamBasic.cc b/catch/performance/stream/hipStreamBasic.cc new file mode 100644 index 0000000000..ffa95b00d7 --- /dev/null +++ b/catch/performance/stream/hipStreamBasic.cc @@ -0,0 +1,269 @@ +/* +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 +#include + +/** + * @addtogroup stream stream + * @{ + * @ingroup PerformanceTest + * Contains performance tests for all hipStream related APIs + */ + + class HipDeviceGetStreamPriorityRangeBenchmark : public Benchmark { + public: + void operator()() { + int priority_min, priority_max; + TIMED_SECTION(kTimerTypeCpu) { HIP_CHECK(hipDeviceGetStreamPriorityRange(&priority_min, &priority_max)); } + } +}; + +class HipStreamQueryBenchmark : public Benchmark { + public: + void operator()(bool perform_work) { + hipError_t error; + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + void *dptr; + + if(perform_work) { + HIP_CHECK(hipMallocAsync(&dptr, 2048 * 4, stream)); + } + + TIMED_SECTION(kTimerTypeCpu) { error = hipStreamQuery(stream); } + + if(perform_work) { + HIP_CHECK(hipFreeAsync(dptr, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + } + + HIP_CHECK(hipStreamDestroy(stream)); + } +}; + +class HipStreamSynchronizeBenchmark : public Benchmark { + public: + void operator()() { + hipError_t error; + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + TIMED_SECTION(kTimerTypeCpu) { error = hipStreamSynchronize(stream); } + + HIP_CHECK(hipStreamDestroy(stream)); + } +}; + +class HipStreamDestroyBenchmark : public Benchmark { + public: + void operator()() { + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + TIMED_SECTION(kTimerTypeCpu) { HIP_CHECK(hipStreamDestroy(stream)); } + } +}; + +class HipStreamCreateBenchmark : public Benchmark { + public: + void operator()() { + hipStream_t stream; + + TIMED_SECTION(kTimerTypeCpu) { HIP_CHECK(hipStreamCreate(&stream)); } + + HIP_CHECK(hipStreamDestroy(stream)); + } +}; + +class HipStreamCreateWithPriorityBenchmark : public Benchmark { + public: + void operator()(unsigned int flag) { + hipStream_t stream; + int priority_min, priority_max, priority_mid; + + HIP_CHECK(hipDeviceGetStreamPriorityRange(&priority_min, &priority_max)); + priority_mid = (priority_max + priority_min) / 2; + + TIMED_SECTION(kTimerTypeCpu) { HIP_CHECK(hipStreamCreateWithPriority(&stream, flag, priority_mid)); } + + HIP_CHECK(hipStreamDestroy(stream)); + } +}; + + + +static std::string GetStreamCreateFlagName(unsigned flag) { + switch (flag) { + case hipStreamDefault: + return "hipStreamDefault"; + case hipStreamNonBlocking: + return "hipStreamNonBlocking"; + default: + return "flag combination"; + } +} + +class HipStreamCreateWithFlagsBenchmark : public Benchmark { + public: + void operator()(unsigned int flag) { + hipStream_t stream; + + TIMED_SECTION(kTimerTypeCpu) { HIP_CHECK(hipStreamCreateWithFlags(&stream, flag)); } + + HIP_CHECK(hipStreamDestroy(stream)); + } +}; + +/** + * Test Description + * ------------------------ + * - Executes `hipStreamCreate`: + * Test source + * ------------------------ + * - performance/stream/hipStreamBasic.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Performance_hipStreamCreate") { + HipStreamCreateBenchmark benchmark; + benchmark.Run(); +} + +static void RunBenchmark(unsigned flag) { + HipStreamCreateWithFlagsBenchmark benchmark; + benchmark.AddSectionName(GetStreamCreateFlagName(flag)); + benchmark.Run(flag); +} + +static void RunBenchmarkWithPriority(unsigned flag) { + HipStreamCreateWithPriorityBenchmark benchmark; + benchmark.AddSectionName(GetStreamCreateFlagName(flag)); + benchmark.Run(flag); +} + +/** + * Test Description + * ------------------------ + * - Executes `hipStreamCreateWithFlags` with all flags: + * -# Flags + * - hipStreamDefault + * - hipStreamNonBlocking + * Test source + * ------------------------ + * - performance/stream/hipStreamBasic.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Performance_hipStreamCreateWithFlags") { + const auto flag = GENERATE(hipStreamDefault, hipStreamNonBlocking); + RunBenchmark(flag); +} + +/** + * Test Description + * ------------------------ + * - Executes `hipStreamCreateWithPriority` with all flags: + * -# Flags + * - hipStreamDefault + * - hipStreamNonBlocking + * Test source + * ------------------------ + * - performance/stream/hipStreamBasic.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Performance_hipStreamCreateWithPriority") { + const auto flag = GENERATE(hipStreamDefault, hipStreamNonBlocking); + RunBenchmarkWithPriority(flag); +} + +/** + * Test Description + * ------------------------ + * - Executes `hipStreamDestroy`: + * Test source + * ------------------------ + * - performance/stream/hipStreamBasic.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Performance_hipStreamDestroy") { + HipStreamDestroyBenchmark benchmark; + benchmark.Run(); +} + +/** + * Test Description + * ------------------------ + * - Executes `hipDeviceGetStreamPriorityRange`: + * Test source + * ------------------------ + * - performance/stream/hipStreamBasic.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Performance_hipDeviceGetStreamPriorityRange") { + HipDeviceGetStreamPriorityRangeBenchmark benchmark; + benchmark.Run(); +} + +/** + * Test Description + * ------------------------ + * - Executes `hipStreamQuery`: + * Test source + * ------------------------ + * - performance/stream/hipStreamBasic.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Performance_hipStreamQuery") { + const auto perform_work = GENERATE(true, false); + HipStreamQueryBenchmark benchmark; + if(perform_work) { + benchmark.AddSectionName("stream with work"); + } else { + benchmark.AddSectionName("stream without work"); + } + benchmark.Run(perform_work); +} + +/** + * Test Description + * ------------------------ + * - Executes `hipDeviceGetStreamPriorityRange`: + * Test source + * ------------------------ + * - performance/stream/hipStreamBasic.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Performance_hipStreamSynchronize") { + HipStreamSynchronizeBenchmark benchmark; + benchmark.Run(); +} diff --git a/catch/performance/stream/hipStreamGetFlags.cc b/catch/performance/stream/hipStreamGetFlags.cc new file mode 100644 index 0000000000..ffead22f39 --- /dev/null +++ b/catch/performance/stream/hipStreamGetFlags.cc @@ -0,0 +1,75 @@ +/* +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 + +/** + * @addtogroup stream stream + * @{ + * @ingroup PerformanceTest + */ + +class StreamGetFlagsBenchmark : public Benchmark { + public: + void operator()(unsigned int expected_flag) { + unsigned int returned_flags{}; + hipStream_t stream; + + HIP_CHECK(hipStreamCreateWithFlags(&stream, expected_flag)); + TIMED_SECTION(kTimerTypeCpu) { + HIP_CHECK(hipStreamGetFlags(stream, &returned_flags)) + } + HIP_CHECK(hipStreamDestroy(stream)); + } +}; + +static void RunBenchmark(unsigned int expected_flag) { + StreamGetFlagsBenchmark benchmark; + switch (expected_flag) { + case hipStreamDefault: + benchmark.AddSectionName("hipStreamDefault"); + break; + case hipStreamNonBlocking: + benchmark.AddSectionName("hipStreamNonBlocking"); + break; + default: + benchmark.AddSectionName("unknown flag type"); + } + benchmark.Run(expected_flag); +} + +/** + * Test Description + * ------------------------ + * - Executes `hipStreamGetFlags`: + * -# Flags: + * - `hipStreamDefault` + * - `hipStreamNonBlocking` + * Test source + * ------------------------ + * - performance/stream/hipStreamGetFlags.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Performance_hipStreamGetFlags") { + unsigned int expected_flag = GENERATE(hipStreamDefault, hipStreamNonBlocking); + RunBenchmark(expected_flag); +} diff --git a/catch/performance/stream/hipStreamGetPriority.cc b/catch/performance/stream/hipStreamGetPriority.cc new file mode 100644 index 0000000000..993471a251 --- /dev/null +++ b/catch/performance/stream/hipStreamGetPriority.cc @@ -0,0 +1,74 @@ +/* +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 + +/** + * @addtogroup stream stream + * @{ + * @ingroup PerformanceTest + */ + +class StreamGetPriorityBenchmark : public Benchmark { + public: + void operator()(Streams stream_type) { + const StreamGuard stream_guard{stream_type}; + const hipStream_t stream = stream_guard.stream(); + + int priority{}; + TIMED_SECTION(kTimerTypeCpu) { + HIP_CHECK(hipStreamGetPriority(stream, &priority)); + } + } +}; + +static void RunBenchmark(Streams stream_type) { + StreamGetPriorityBenchmark benchmark; + switch (stream_type) { + case Streams::nullstream: + benchmark.AddSectionName("null stream"); + break; + case Streams::created: + benchmark.AddSectionName("created"); + break; + default: + benchmark.AddSectionName("per thread stream"); + } + benchmark.Run(stream_type); +} + +/** + * Test Description + * ------------------------ + * - Executes `hipStreamGetPriority`: + * -# Stream types: + * - `null` + * - created + * Test source + * ------------------------ + * - performance/stream/hipStreamGetPriority.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Performance_hipStreamGetPriority") { + Streams stream_type = GENERATE(Streams::nullstream, Streams::created); + RunBenchmark(stream_type); +} \ No newline at end of file diff --git a/catch/performance/stream/hipStreamWaitEvent.cc b/catch/performance/stream/hipStreamWaitEvent.cc new file mode 100644 index 0000000000..6810b5b619 --- /dev/null +++ b/catch/performance/stream/hipStreamWaitEvent.cc @@ -0,0 +1,80 @@ +/* +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 + +/** + * @addtogroup stream stream + * @{ + * @ingroup PerformanceTest + */ + +class StreamWaitEventBenchmark : public Benchmark { + public: + void operator()(Streams stream_type) { + const StreamGuard stream_guard{stream_type}; + const hipStream_t stream = stream_guard.stream(); + hipEvent_t wait_event{nullptr}; + + HIP_CHECK(hipEventCreate(&wait_event)); + REQUIRE(wait_event != nullptr); + HIP_CHECK(hipEventRecord(wait_event, stream)); + + TIMED_SECTION(kTimerTypeCpu) { + HIP_CHECK(hipStreamWaitEvent(stream, wait_event, 0)); + HIP_CHECK(hipStreamSynchronize(stream)); + } + HIP_CHECK(hipEventDestroy(wait_event)); + } +}; + +static void RunBenchmark(Streams stream_type) { + StreamWaitEventBenchmark benchmark{}; + switch (stream_type) { + case Streams::nullstream: + benchmark.AddSectionName("null stream"); + break; + case Streams::created: + benchmark.AddSectionName("created"); + break; + default: + benchmark.AddSectionName("per thread stream"); + } + benchmark.Run(stream_type); +} + +/** + * Test Description + * ------------------------ + * - Executes `hipStreamWaitEvent`: + * -# Stream types: + * - `null` + * - created + * Test source + * ------------------------ + * - performance/stream/hipStreamWaitEvent.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Performance_hipStreamWaitEvent") { + Streams stream_type = GENERATE(Streams::nullstream, Streams::created); + RunBenchmark(stream_type); +} diff --git a/catch/performance/stream/hipStreamWaitValue.cc b/catch/performance/stream/hipStreamWaitValue.cc new file mode 100644 index 0000000000..b4c33af2e9 --- /dev/null +++ b/catch/performance/stream/hipStreamWaitValue.cc @@ -0,0 +1,171 @@ +/* +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 + +/** + * @addtogroup stream stream + * @{ + * @ingroup PerformanceTest + */ + +static int IsStreamWaitValueSupported(int device_id) { + int wait_value_supported = 0; +#if HT_AMD + HIP_CHECK(hipDeviceGetAttribute(&wait_value_supported, hipDeviceAttributeCanUseStreamWaitValue, + device_id)); +#else + cuDeviceGetAttribute(&wait_value_supported, CU_DEVICE_ATTRIBUTE_CAN_USE_64_BIT_STREAM_MEM_OPS, + device_id); +#endif + return wait_value_supported; +} + +class StreamWaitValue32Benchmark : public Benchmark { + public: + void operator()(const size_t array_size, unsigned int flag) { + const StreamGuard stream_guard{Streams::created}; + const hipStream_t stream = stream_guard.stream(); + uint32_t* value_ptr; + uint32_t value{0}; + if (flag == hipStreamWaitValueAnd) { + value = 1; + } + HIP_CHECK(hipMalloc(&value_ptr, sizeof(uint32_t) * array_size)); + HIP_CHECK(hipMemset(value_ptr, value, sizeof(uint32_t) * array_size)); + + TIMED_SECTION(kTimerTypeCpu) { + HIP_CHECK(hipStreamWaitValue32(stream, value_ptr, value, flag)); + } + HIP_CHECK(hipFree(value_ptr)); + } +}; + +class StreamWaitValue64Benchmark : public Benchmark { + public: + void operator()(const size_t array_size, unsigned int flag) { + const StreamGuard stream_guard{Streams::created}; + const hipStream_t stream = stream_guard.stream(); + uint64_t* value_ptr; + uint64_t value{0}; + if (flag == hipStreamWaitValueAnd) { + value = 1; + } + HIP_CHECK(hipMalloc(&value_ptr, sizeof(uint64_t) * array_size)); + HIP_CHECK(hipMemset(value_ptr, value, sizeof(uint64_t) * array_size)); + + TIMED_SECTION(kTimerTypeCpu) { + HIP_CHECK(hipStreamWaitValue64(stream, value_ptr, value, flag)); + } + HIP_CHECK(hipFree(value_ptr)); + } +}; + +template +static void RunBenchmark(const size_t array_size, unsigned int flag) { + WaitValueBenchmark benchmark; + benchmark.AddSectionName(std::to_string(array_size)); + switch (flag) { + case hipStreamWaitValueGte: + benchmark.AddSectionName("greater than or equal"); + break; + case hipStreamWaitValueEq: + benchmark.AddSectionName("equal"); + break; + case hipStreamWaitValueAnd: + benchmark.AddSectionName("logical and"); + break; + case hipStreamWaitValueNor: + benchmark.AddSectionName("logical nor"); + break; + default: + benchmark.AddSectionName("unknown flag"); + } + benchmark.Run(array_size, flag); +} + +/** + * Test Description + * ------------------------ + * - Executes `hipStreamWaitValue32` for different array sizes: + * -# 4 KB + * -# 4 MB + * -# 16 MB + * - Uses different flag types for wait criteria: + * -# Greater than or equal + * -# Equal + * -# Logical AND + * -# Logical OR + * Test source + * ------------------------ + * - performance/stream/hipStreamWaitValue.cc + * Test requirements + * ------------------------ + * - Device supports Stream Wait Value operations + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Performance_hipStreamWaitValue32") { +#if HT_AMD + if (!IsStreamWaitValueSupported(0)) { + HipTest::HIP_SKIP_TEST( + "GPU 0 doesn't support hipStreamWaitValue32() function. " + "Hence skipping the testing with Pass result.\n"); + return; + } +#endif + size_t array_size = GENERATE(4_KB, 4_MB, 16_MB); + unsigned int flag = GENERATE(hipStreamWaitValueGte, hipStreamWaitValueEq, hipStreamWaitValueAnd, + hipStreamWaitValueNor); + RunBenchmark(array_size, flag); +} + +/** + * Test Description + * ------------------------ + * - Executes `hipStreamWaitValue64`: + * -# Allocation size: + * - 4 KB + * - 4 MB + * - 16 MB + * -# Wait type: + * - Greater than or equal + * - Equal + * - Logical AND + * - Logical OR + * Test source + * ------------------------ + * - performance/stream/hipStreamWaitValue.cc + * Test requirements + * ------------------------ + * - Device supports Stream Wait Value operations + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Performance_hipStreamWaitValue64") { + if (!IsStreamWaitValueSupported(0)) { + HipTest::HIP_SKIP_TEST( + "GPU 0 doesn't support hipStreamWaitValue64() function. " + "Hence skipping the testing with Pass result.\n"); + return; + } + size_t array_size = GENERATE(4_KB, 4_MB, 16_MB); + unsigned int flag = GENERATE(hipStreamWaitValueGte, hipStreamWaitValueEq, hipStreamWaitValueAnd, + hipStreamWaitValueNor); + RunBenchmark(array_size, flag); +} diff --git a/catch/performance/stream/hipStreamWriteValue.cc b/catch/performance/stream/hipStreamWriteValue.cc new file mode 100644 index 0000000000..56c94393bc --- /dev/null +++ b/catch/performance/stream/hipStreamWriteValue.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 +#include + +/** + * @addtogroup stream stream + * @{ + * @ingroup PerformanceTest + */ + +#if HT_NVIDIA +static int IsStreamWriteValueSupported(int device_id) { + int write_value_supported = 0; + + cuDeviceGetAttribute(&write_value_supported, CU_DEVICE_ATTRIBUTE_CAN_USE_64_BIT_STREAM_MEM_OPS, + device_id); + return write_value_supported; +} +#endif + +class StreamWriteValue32Benchmark : public Benchmark { + public: + void operator()(const size_t array_size) { + const StreamGuard stream_guard{Streams::created}; + const hipStream_t stream = stream_guard.stream(); + uint32_t* value_ptr; + uint32_t value{0}; + HIP_CHECK(hipMalloc(&value_ptr, sizeof(uint32_t) * array_size)); + HIP_CHECK(hipMemset(value_ptr, value, sizeof(uint32_t) * array_size)); + + TIMED_SECTION(kTimerTypeCpu) { HIP_CHECK(hipStreamWriteValue32(stream, value_ptr, value, 0)); } + HIP_CHECK(hipFree(value_ptr)); + } +}; + +class StreamWriteValue64Benchmark : public Benchmark { + public: + void operator()(const size_t array_size) { + const StreamGuard stream_guard{Streams::created}; + const hipStream_t stream = stream_guard.stream(); + uint64_t* value_ptr; + uint64_t value{0}; + HIP_CHECK(hipMalloc(&value_ptr, sizeof(uint64_t) * array_size)); + HIP_CHECK(hipMemset(value_ptr, value, sizeof(uint64_t) * array_size)); + + TIMED_SECTION(kTimerTypeCpu) { HIP_CHECK(hipStreamWriteValue64(stream, value_ptr, value, 0)); } + HIP_CHECK(hipFree(value_ptr)); + } +}; + +template static void RunBenchmark(const size_t array_size) { + WriteValueBenchmark benchmark; + benchmark.AddSectionName(std::to_string(array_size)); + benchmark.Run(array_size); +} + +/** + * Test Description + * ------------------------ + * - Executes `hipStreamWriteValue32`: + * -# Allocation size: + * - 4 KB + * - 4 MB + * - 16 MB + * Test source + * ------------------------ + * - performance/stream/hipStreamWriteValue.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Performance_hipStreamWriteValue32") { + size_t array_size = GENERATE(4_KB, 4_MB, 16_MB); + RunBenchmark(array_size); +} + +/** + * Test Description + * ------------------------ + * - Executes `hipStreamWriteValue64`: + * -# Allocation size: + * - 4 KB + * - 4 MB + * - 16 MB + * Test source + * ------------------------ + * - performance/stream/hipStreamWriteValue.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Performance_hipStreamWriteValue64") { +#if HT_NVIDIA + if (!IsStreamWriteValueSupported(0)) { + HipTest::HIP_SKIP_TEST( + "GPU 0 doesn't support hipStreamWriteValue64() function. " + "Hence skipping the testing with Pass result.\n"); + return; + } +#endif + size_t array_size = GENERATE(4_KB, 4_MB, 16_MB); + RunBenchmark(array_size); +} diff --git a/catch/performance/stream/mem_pools_performance_common.hh b/catch/performance/stream/mem_pools_performance_common.hh new file mode 100644 index 0000000000..b6f6f4b489 --- /dev/null +++ b/catch/performance/stream/mem_pools_performance_common.hh @@ -0,0 +1,74 @@ +/* +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 + +#if __linux__ + static const hipMemAllocationHandleType kHandleType = hipMemHandleTypePosixFileDescriptor; +#else + static const hipMemAllocationHandleType kHandleType = hipMemHandleTypeWin32; +#endif + +static int AreMemPoolsSupported(int device_id) { + int mem_pools_supported = 0; + HIP_CHECK(hipDeviceGetAttribute(&mem_pools_supported, + hipDeviceAttributeMemoryPoolsSupported, 0)); + return mem_pools_supported; +} + +static hipMemPoolProps CreateMemPoolProps(const int device_id, const hipMemAllocationHandleType handle_type) { + hipMemPoolProps kPoolProps = { + hipMemAllocationTypePinned, + handle_type, + { + hipMemLocationTypeDevice, + device_id + }, + nullptr, + {0} + }; + + return kPoolProps; +} + +static std::string GetMemPoolAttrSectionName(const hipMemPoolAttr attribute) { + switch (attribute) { + case hipMemPoolReuseFollowEventDependencies: + return "ReuseFollowEventDependencies"; + case hipMemPoolReuseAllowOpportunistic: + return "ReuseAllowOpportunistic"; + case hipMemPoolReuseAllowInternalDependencies: + return "ReuseAllowInternalDependencies"; + case hipMemPoolAttrReleaseThreshold: + return "AttrReleaseThreshold"; + case hipMemPoolAttrReservedMemCurrent: + return "AttrReservedMemCurrent"; + case hipMemPoolAttrReservedMemHigh: + return "AttrReservedMemHigh"; + case hipMemPoolAttrUsedMemCurrent: + return "AttrUsedMemCurrent"; + case hipMemPoolAttrUsedMemHigh: + return "AttrUsedMemHigh"; + default: + return "unknown attribute"; + } +}