From 101abc7b392b4f9279e4f0fd50123aa82f9a3e66 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Mon, 14 Aug 2023 21:17:55 +0530 Subject: [PATCH] SWDEV-403471 - [catch2][dtest] Converting perftests-memory files from HIT to catch2 (#342) Change-Id: I13d2513f31dffe0b280039c888a97cc0d7bba31f [ROCm/hip-tests commit: cf174d5a470946c9c836a3e7394b47b50e826431] --- projects/hip-tests/catch/CMakeLists.txt | 1 + .../hip-tests/catch/perftests/CMakeLists.txt | 24 + .../catch/perftests/memory/CMakeLists.txt | 47 ++ .../memory/hipPerfBufferCopyRectSpeed.cc | 242 ++++++++ .../memory/hipPerfBufferCopySpeed.cc | 258 +++++++++ .../memory/hipPerfDevMemReadSpeed.cc | 154 +++++ .../memory/hipPerfDevMemWriteSpeed.cc | 147 +++++ .../perftests/memory/hipPerfHostNumaAlloc.cc | 191 ++++++ .../catch/perftests/memory/hipPerfMemFill.cc | 545 ++++++++++++++++++ .../memory/hipPerfMemMallocCpyFree.cc | 144 +++++ .../catch/perftests/memory/hipPerfMemcpy.cc | 118 ++++ .../catch/perftests/memory/hipPerfMemset.cc | 419 ++++++++++++++ .../perftests/memory/hipPerfSampleRate.cc | 311 ++++++++++ .../memory/hipPerfSharedMemReadSpeed.cc | 262 +++++++++ 14 files changed, 2863 insertions(+) create mode 100644 projects/hip-tests/catch/perftests/CMakeLists.txt create mode 100644 projects/hip-tests/catch/perftests/memory/CMakeLists.txt create mode 100644 projects/hip-tests/catch/perftests/memory/hipPerfBufferCopyRectSpeed.cc create mode 100644 projects/hip-tests/catch/perftests/memory/hipPerfBufferCopySpeed.cc create mode 100644 projects/hip-tests/catch/perftests/memory/hipPerfDevMemReadSpeed.cc create mode 100644 projects/hip-tests/catch/perftests/memory/hipPerfDevMemWriteSpeed.cc create mode 100644 projects/hip-tests/catch/perftests/memory/hipPerfHostNumaAlloc.cc create mode 100644 projects/hip-tests/catch/perftests/memory/hipPerfMemFill.cc create mode 100644 projects/hip-tests/catch/perftests/memory/hipPerfMemMallocCpyFree.cc create mode 100644 projects/hip-tests/catch/perftests/memory/hipPerfMemcpy.cc create mode 100644 projects/hip-tests/catch/perftests/memory/hipPerfMemset.cc create mode 100644 projects/hip-tests/catch/perftests/memory/hipPerfSampleRate.cc create mode 100644 projects/hip-tests/catch/perftests/memory/hipPerfSharedMemReadSpeed.cc diff --git a/projects/hip-tests/catch/CMakeLists.txt b/projects/hip-tests/catch/CMakeLists.txt index e34f59180d..8f7aa92624 100644 --- a/projects/hip-tests/catch/CMakeLists.txt +++ b/projects/hip-tests/catch/CMakeLists.txt @@ -241,6 +241,7 @@ add_subdirectory(kernels ${CATCH_BUILD_DIR}/kernels) add_subdirectory(hipTestMain ${CATCH_BUILD_DIR}/hipTestMain) add_subdirectory(stress ${CATCH_BUILD_DIR}/stress) add_subdirectory(TypeQualifiers ${CATCH_BUILD_DIR}/TypeQualifiers) +add_subdirectory(perftests ${CATCH_BUILD_DIR}/perftests) if(UNIX) add_subdirectory(multiproc ${CATCH_BUILD_DIR}/multiproc) endif() diff --git a/projects/hip-tests/catch/perftests/CMakeLists.txt b/projects/hip-tests/catch/perftests/CMakeLists.txt new file mode 100644 index 0000000000..be81c63a12 --- /dev/null +++ b/projects/hip-tests/catch/perftests/CMakeLists.txt @@ -0,0 +1,24 @@ +# Copyright (c) 2023 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. + +add_custom_target(perf_test COMMAND "${CMAKE_CTEST_COMMAND}" -R "Perf_" + COMMENT "Build complete, now executing the performnce test ...") + +add_subdirectory(memory) diff --git a/projects/hip-tests/catch/perftests/memory/CMakeLists.txt b/projects/hip-tests/catch/perftests/memory/CMakeLists.txt new file mode 100644 index 0000000000..542fdfe3f8 --- /dev/null +++ b/projects/hip-tests/catch/perftests/memory/CMakeLists.txt @@ -0,0 +1,47 @@ +# Copyright (c) 2023 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. + +set(TEST_SRC + hipPerfMemcpy.cc + hipPerfBufferCopyRectSpeed.cc + hipPerfBufferCopySpeed.cc + hipPerfDevMemReadSpeed.cc + hipPerfDevMemWriteSpeed.cc + hipPerfMemFill.cc + hipPerfMemMallocCpyFree.cc + hipPerfMemset.cc + hipPerfSampleRate.cc + hipPerfSharedMemReadSpeed.cc +) + +if(HIP_PLATFORM MATCHES "amd") + set(TEST_SRC + ${TEST_SRC} + hipPerfHostNumaAlloc.cc) + +hip_add_exe_to_target(NAME perfMemoryTest + TEST_SRC ${TEST_SRC} + TEST_TARGET_NAME perf_test + LINKER_LIBS numa) +else() +hip_add_exe_to_target(NAME perfMemoryTest + TEST_SRC ${TEST_SRC} + TEST_TARGET_NAME perf_test) +endif() diff --git a/projects/hip-tests/catch/perftests/memory/hipPerfBufferCopyRectSpeed.cc b/projects/hip-tests/catch/perftests/memory/hipPerfBufferCopyRectSpeed.cc new file mode 100644 index 0000000000..2676140a15 --- /dev/null +++ b/projects/hip-tests/catch/perftests/memory/hipPerfBufferCopyRectSpeed.cc @@ -0,0 +1,242 @@ +/* +Copyright (c) 2023 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. +*/ + +/** +* @addtogroup hipMemcpy2DAsync hipMemcpy2DAsync +* @{ +* @ingroup perfMemoryTest +* `hipMemcpy2DAsync(void* dst, size_t dpitch, const void* src, size_t spitch, +* size_t width, size_t height, hipMemcpyKind kind, hipStream_t stream = 0)` - +* Copies data between host and device. +*/ + +#include + +#define NUM_SIZES 8 +// 4KB, 8KB, 64KB, 256KB, 1 MB, 4MB, 16 MB, 16MB+10 +static const unsigned int Sizes[NUM_SIZES] = + {4096, 8192, 65536, 262144, 1048576, 4194304, 16777216, 16777216+10}; + +static const unsigned int Iterations[2] = {1, 1000}; + +#define BUF_TYPES 4 +// 16 ways to combine 4 different buffer types +#define NUM_SUBTESTS (BUF_TYPES*BUF_TYPES) + +static void setData(void *ptr, unsigned int size, char value) { + char *ptr2 = reinterpret_cast(ptr); + for (unsigned int i = 0; i < size ; i++) { + ptr2[i] = value; + } +} + +static bool hipPerfBufferCopyRectSpeed_test(int p_tests) { + unsigned int bufSize_; + unsigned int numIter; + bool hostMalloc[2] = {false}; + bool hostRegister[2] = {false}; + bool unpinnedMalloc[2] = {false}; + void *memptr[2] = {NULL}; + void *alignedmemptr[2] = {NULL}; + void *srcBuffer = NULL; + void *dstBuffer = NULL; + + int numTests = (p_tests == -1) ? (NUM_SIZES*NUM_SUBTESTS*2 - 1) : p_tests; + int test = (p_tests == -1) ? 0 : p_tests; + + for ( ; test <= numTests ; test++ ) { + unsigned int srcTest = (test / NUM_SIZES) % BUF_TYPES; + unsigned int dstTest = (test / (NUM_SIZES*BUF_TYPES)) % BUF_TYPES; + bufSize_ = Sizes[test % NUM_SIZES]; + hostMalloc[0] = hostMalloc[1] = false; + hostRegister[0] = hostRegister[1] = false; + unpinnedMalloc[0] = unpinnedMalloc[1] = false; + srcBuffer = dstBuffer = 0; + memptr[0] = memptr[1] = 0; + alignedmemptr[0] = alignedmemptr[1] = NULL; + + size_t width = static_cast(sqrt(static_cast(bufSize_))); + + if (srcTest == 3) { + hostRegister[0] = true; + } else if (srcTest == 2) { + hostMalloc[0] = true; + } else if (srcTest == 1) { + unpinnedMalloc[0] = true; + } + + if (dstTest == 1) { + unpinnedMalloc[1] = true; + } else if (dstTest == 2) { + hostMalloc[1] = true; + } else if (dstTest == 3) { + hostRegister[1] = true; + } + + numIter = Iterations[test / (NUM_SIZES * NUM_SUBTESTS)]; + + if (hostMalloc[0]) { + HIP_CHECK(hipHostMalloc(reinterpret_cast(&srcBuffer), + bufSize_, 0)); + setData(srcBuffer, bufSize_, 0xd0); + } else if (hostRegister[0]) { + memptr[0] = malloc(bufSize_ + 4096); + alignedmemptr[0] = reinterpret_cast(memptr[0]); + srcBuffer = alignedmemptr[0]; + setData(srcBuffer, bufSize_, 0xd0); + HIP_CHECK(hipHostRegister(srcBuffer, bufSize_, 0)); + } else if (unpinnedMalloc[0]) { + memptr[0] = malloc(bufSize_ + 4096); + alignedmemptr[0] = reinterpret_cast(memptr[0]); + srcBuffer = alignedmemptr[0]; + setData(srcBuffer, bufSize_, 0xd0); + } else { + HIP_CHECK(hipMalloc(&srcBuffer, bufSize_)); + HIP_CHECK(hipMemset(srcBuffer, 0xd0, bufSize_)); + } + + if (hostMalloc[1]) { + HIP_CHECK(hipHostMalloc(reinterpret_cast(&dstBuffer), + bufSize_, 0)); + } else if (hostRegister[1]) { + memptr[1] = malloc(bufSize_ + 4096); + alignedmemptr[1] = reinterpret_cast(memptr[0]); + dstBuffer = alignedmemptr[1]; + HIP_CHECK(hipHostRegister(dstBuffer, bufSize_, 0)); + } else if (unpinnedMalloc[1]) { + memptr[1] = malloc(bufSize_ + 4096); + alignedmemptr[1] = reinterpret_cast(memptr[0]); + dstBuffer = alignedmemptr[1]; + } else { + HIP_CHECK(hipMalloc(&dstBuffer, bufSize_)); + } + + // warm up + HIP_CHECK(hipMemcpy2D(dstBuffer, width, srcBuffer, + width, width, width, hipMemcpyDefault)); + + // measure performance based on host time + auto all_start = std::chrono::steady_clock::now(); + + for (unsigned int i = 0; i < numIter; i++) { + HIP_CHECK(hipMemcpy2DAsync(dstBuffer, width, srcBuffer, + width, width, width, hipMemcpyDefault, NULL)); + } + HIP_CHECK(hipDeviceSynchronize()); + + auto all_end = std::chrono::steady_clock::now(); + std::chrono::duration elapsed_secs = all_end - all_start; + + // read speed in GB/s + double perf = (static_cast(bufSize_ * numIter) * + static_cast(1e-09)) / elapsed_secs.count(); + + const char *strSrc = NULL; + const char *strDst = NULL; + if (hostMalloc[0]) + strSrc = "hHM"; + else if (hostRegister[0]) + strSrc = "hHR"; + else if (unpinnedMalloc[0]) + strSrc = "unp"; + else + strSrc = "hM"; + + if (hostMalloc[1]) + strDst = "hHM"; + else if (hostRegister[1]) + strDst = "hHR"; + else if (unpinnedMalloc[1]) + strDst = "unp"; + else + strDst = "hM"; + + // Double results when src and dst are both on device + if ((!hostMalloc[0] && !hostRegister[0] && !unpinnedMalloc[0]) && + (!hostMalloc[1] && !hostRegister[1] && !unpinnedMalloc[1])) + perf *= 2.0; + // Double results when src and dst are both in sysmem + if ((hostMalloc[0] || hostRegister[0] || unpinnedMalloc[0]) && + (hostMalloc[1] || hostRegister[1] || unpinnedMalloc[1])) + perf *= 2.0; + + INFO("hipPerfBufferCopyRectSpeed[" << test << "]\t( " << bufSize_ << + ")\ts:" << strSrc << " d:" << strDst << "\ti:" << numIter << + "\t(GB/s) perf\t" << (float)perf); + + // Free src + if (hostMalloc[0]) { + HIP_CHECK(hipHostFree(srcBuffer)); + } else if (hostRegister[0]) { + HIP_CHECK(hipHostUnregister(srcBuffer)); + free(memptr[0]); + } else if (unpinnedMalloc[0]) { + free(memptr[0]); + } else { + HIP_CHECK(hipFree(srcBuffer)); + } + + // Free dst + if (hostMalloc[1]) { + HIP_CHECK(hipHostFree(dstBuffer)); + } else if (hostRegister[1]) { + HIP_CHECK(hipHostUnregister(dstBuffer)); + free(memptr[1]); + } else if (unpinnedMalloc[1]) { + free(memptr[1]); + } else { + HIP_CHECK(hipFree(dstBuffer)); + } + } + return true; +} + +/** +* Test Description +* ------------------------ +*  - Verify hipPerfBufferCopy status. +* Test source +* ------------------------ +*  - perftests/memory/hipPerfBufferCopyRectSpeed.cc +* Test requirements +* ------------------------ +*  - HIP_VERSION >= 5.6 +*/ + +TEST_CASE("Perf_hipPerfBufferCopyRectSpeed_test") { + int numDevices = 0; + HIP_CHECK(hipGetDeviceCount(&numDevices)); + + if (numDevices <= 0) { + SUCCEED("Skipped testcase hipPerfBufferCopyRectSpeed" + "as there is no device to test."); + } else { + int deviceId = 0; + HIP_CHECK(hipSetDevice(deviceId)); + hipDeviceProp_t props = {0}; + HIP_CHECK(hipGetDeviceProperties(&props, deviceId)); + + INFO("hipPerfBufferCopyRectSpeed - info: Set device to " << deviceId + << " : " << props.name << "Legend: unp - unpinned(malloc)," + " hM - hipMalloc(device)\n hHR - hipHostRegister(pinned)," + " hHM - hipHostMalloc(prePinned)\n"); + + REQUIRE(true == hipPerfBufferCopyRectSpeed_test(1)); + } +} diff --git a/projects/hip-tests/catch/perftests/memory/hipPerfBufferCopySpeed.cc b/projects/hip-tests/catch/perftests/memory/hipPerfBufferCopySpeed.cc new file mode 100644 index 0000000000..f9201bd672 --- /dev/null +++ b/projects/hip-tests/catch/perftests/memory/hipPerfBufferCopySpeed.cc @@ -0,0 +1,258 @@ +/* +Copyright (c) 2023 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. +*/ + +/** +* @addtogroup hipMemcpyAsync hipMemcpyAsync +* @{ +* @ingroup perfMemoryTest +* `hipMemcpyAsync(void* dst, const void* src, size_t count, +* hipMemcpyKind kind, hipStream_t stream = 0)` - +* Copies data between host and device. +*/ + +#include + +#define NUM_SIZES 9 +// 4KB, 8KB, 64KB, 256KB, 1 MB, 4MB, 16 MB, 16MB+10 +static const unsigned int Sizes[NUM_SIZES] = + {4096, 8192, 65536, 262144, 524288, 1048576, 4194304, 16777216, 16777216+10}; + +static const unsigned int Iterations[2] = {1, 1000}; + +#define BUF_TYPES 4 +// 16 ways to combine 4 different buffer types +#define NUM_SUBTESTS (BUF_TYPES*BUF_TYPES) + +static void setData(void *ptr, unsigned int size, char value) { + char *ptr2 = reinterpret_cast(ptr); + for (unsigned int i = 0; i < size ; i++) { + ptr2[i] = value; + } +} + +static void checkData(void *ptr, unsigned int size, char value) { + char *ptr2 = reinterpret_cast(ptr); + for (unsigned int i = 0; i < size; i++) { + if (ptr2[i] != value) { + INFO("Validation failed at " << i << " Got " << ptr2[i] << + " Expected " << value); + REQUIRE(false); + } + } +} + +static bool hipPerfBufferCopySpeed_test(int p_tests) { + unsigned int bufSize_; + unsigned int numIter; + bool hostMalloc[2] = {false}; + bool hostRegister[2] = {false}; + bool unpinnedMalloc[2] = {false}; + void *memptr[2] = {NULL}; + void *alignedmemptr[2] = {NULL}; + void *srcBuffer = NULL; + void *dstBuffer = NULL; + + int numTests = (p_tests == -1) ? (NUM_SIZES*NUM_SUBTESTS*2 - 1) : p_tests; + int test = (p_tests == -1) ? 0 : p_tests; + + for ( ; test <= numTests; test++ ) { + unsigned int srcTest = (test / NUM_SIZES) % BUF_TYPES; + unsigned int dstTest = (test / (NUM_SIZES*BUF_TYPES)) % BUF_TYPES; + bufSize_ = Sizes[test % NUM_SIZES]; + hostMalloc[0] = hostMalloc[1] = false; + hostRegister[0] = hostRegister[1] = false; + unpinnedMalloc[0] = unpinnedMalloc[1] = false; + srcBuffer = dstBuffer = 0; + memptr[0] = memptr[1] = NULL; + alignedmemptr[0] = alignedmemptr[1] = NULL; + + if (srcTest == 3) { + hostRegister[0] = true; + } else if (srcTest == 2) { + hostMalloc[0] = true; + } else if (srcTest == 1) { + unpinnedMalloc[0] = true; + } + + if (dstTest == 1) { + unpinnedMalloc[1] = true; + } else if (dstTest == 2) { + hostMalloc[1] = true; + } else if (dstTest == 3) { + hostRegister[1] = true; + } + + numIter = Iterations[test / (NUM_SIZES * NUM_SUBTESTS)]; + + if (hostMalloc[0]) { + HIP_CHECK(hipHostMalloc(reinterpret_cast(&srcBuffer), + bufSize_, 0)); + setData(srcBuffer, bufSize_, 0xd0); + } else if (hostRegister[0]) { + memptr[0] = malloc(bufSize_ + 4096); + alignedmemptr[0] = reinterpret_cast(memptr[0]); + srcBuffer = alignedmemptr[0]; + setData(srcBuffer, bufSize_, 0xd0); + HIP_CHECK(hipHostRegister(srcBuffer, bufSize_, 0)); + } else if (unpinnedMalloc[0]) { + memptr[0] = malloc(bufSize_ + 4096); + alignedmemptr[0] = reinterpret_cast(memptr[0]); + srcBuffer = alignedmemptr[0]; + setData(srcBuffer, bufSize_, 0xd0); + } else { + HIP_CHECK(hipMalloc(&srcBuffer, bufSize_)); + HIP_CHECK(hipMemset(srcBuffer, 0xd0, bufSize_)); + } + + if (hostMalloc[1]) { + HIP_CHECK(hipHostMalloc(reinterpret_cast(&dstBuffer), + bufSize_, 0)); + } else if (hostRegister[1]) { + memptr[1] = malloc(bufSize_ + 4096); + alignedmemptr[1] = reinterpret_cast(memptr[1]); + dstBuffer = alignedmemptr[1]; + HIP_CHECK(hipHostRegister(dstBuffer, bufSize_, 0)); + } else if (unpinnedMalloc[1]) { + memptr[1] = malloc(bufSize_ + 4096); + alignedmemptr[1] = reinterpret_cast(memptr[1]); + dstBuffer = alignedmemptr[1]; + } else { + HIP_CHECK(hipMalloc(&dstBuffer, bufSize_)); + } + + // warm up + HIP_CHECK(hipMemcpy(dstBuffer, srcBuffer, bufSize_, hipMemcpyDefault)); + + // measure performance based on host time + auto all_start = std::chrono::steady_clock::now(); + + for (unsigned int i = 0; i < numIter; i++) { + HIP_CHECK(hipMemcpyAsync(dstBuffer, srcBuffer, bufSize_, + hipMemcpyDefault, NULL)); + } + HIP_CHECK(hipDeviceSynchronize()); + + auto all_end = std::chrono::steady_clock::now(); + std::chrono::duration elapsed_secs = all_end - all_start; + + // read speed in GB/s + double perf = (static_cast(bufSize_ * numIter) * + static_cast(1e-09)) / elapsed_secs.count(); + + const char *strSrc = NULL; + const char *strDst = NULL; + if (hostMalloc[0]) + strSrc = "hHM"; + else if (hostRegister[0]) + strSrc = "hHR"; + else if (unpinnedMalloc[0]) + strSrc = "unp"; + else + strSrc = "hM"; + + if (hostMalloc[1]) + strDst = "hHM"; + else if (hostRegister[1]) + strDst = "hHR"; + else if (unpinnedMalloc[1]) + strDst = "unp"; + else + strDst = "hM"; + + // Double results when src and dst are both on device + if ((!hostMalloc[0] && !hostRegister[0] && !unpinnedMalloc[0]) && + (!hostMalloc[1] && !hostRegister[1] && !unpinnedMalloc[1])) + perf *= 2.0; + // Double results when src and dst are both in sysmem + if ((hostMalloc[0] || hostRegister[0] || unpinnedMalloc[0]) && + (hostMalloc[1] || hostRegister[1] || unpinnedMalloc[1])) + perf *= 2.0; + + INFO("HIPPerfBufferCopySpeed[" << test << "]\t( " << bufSize_ << + ")\ts:" << strSrc << " d:" << strDst << "\ti:" << numIter << + "\t(GB/s) perf\t" << (float)perf); + + // Verification + void* temp = malloc(bufSize_ + 4096); + void* chkBuf = reinterpret_cast(temp); + HIP_CHECK(hipMemcpy(chkBuf, dstBuffer, bufSize_, hipMemcpyDefault)); + checkData(chkBuf, bufSize_, 0xd0); + free(temp); + + // Free src + if (hostMalloc[0]) { + HIP_CHECK(hipHostFree(srcBuffer)); + } else if (hostRegister[0]) { + HIP_CHECK(hipHostUnregister(srcBuffer)); + free(memptr[0]); + } else if (unpinnedMalloc[0]) { + free(memptr[0]); + } else { + HIP_CHECK(hipFree(srcBuffer)); + } + + // Free dst + if (hostMalloc[1]) { + HIP_CHECK(hipHostFree(dstBuffer)); + } else if (hostRegister[1]) { + HIP_CHECK(hipHostUnregister(dstBuffer)); + free(memptr[1]); + } else if (unpinnedMalloc[1]) { + free(memptr[1]); + } else { + HIP_CHECK(hipFree(dstBuffer)); + } + } + + return true; +} + +/** +* Test Description +* ------------------------ +*  - Verify hipPerfBufferCopySpeed status. +* Test source +* ------------------------ +*  - perftests/memory/hipPerfBufferCopySpeed.cc +* Test requirements +* ------------------------ +*  - HIP_VERSION >= 5.6 +*/ + +TEST_CASE("Perf_hipPerfBufferCopySpeed_test") { + int numDevices = 0; + HIP_CHECK(hipGetDeviceCount(&numDevices)); + + if (numDevices <= 0) { + SUCCEED("Skipped testcase hipPerfBufferCopySpeed as" + "there is no device to test."); + } else { + int deviceId = 0; + HIP_CHECK(hipSetDevice(deviceId)); + hipDeviceProp_t props = {0}; + HIP_CHECK(hipGetDeviceProperties(&props, deviceId)); + + INFO("hipPerfBufferCopySpeed - info: Set device to " << deviceId + << " : " << props.name << "Legend: unp - unpinned(malloc)," + " hM - hipMalloc(device)\n hHR - hipHostRegister(pinned)," + " hHM - hipHostMalloc(prePinned)\n"); + + REQUIRE(true == hipPerfBufferCopySpeed_test(1)); + } +} diff --git a/projects/hip-tests/catch/perftests/memory/hipPerfDevMemReadSpeed.cc b/projects/hip-tests/catch/perftests/memory/hipPerfDevMemReadSpeed.cc new file mode 100644 index 0000000000..dad9bfcef4 --- /dev/null +++ b/projects/hip-tests/catch/perftests/memory/hipPerfDevMemReadSpeed.cc @@ -0,0 +1,154 @@ +/* +Copyright (c) 2023 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. +*/ + +/** +* @addtogroup hipMemcpyKernel hipMemcpyKernel +* @{ +* @ingroup perfMemoryTest +* `hipMemcpy(void* dst, const void* src, size_t count, hipMemcpyKind kind)` - +* Copies data between host and device. +*/ + +#include + +#define ARRAY_SIZE 16 + +typedef struct d_uint16 { + uint data[ARRAY_SIZE]; +} d_uint16; + +__global__ static void read_kernel(d_uint16 *src, ulong N, uint *dst) { + size_t idx = (blockIdx.x * blockDim.x + threadIdx.x); + size_t stride = blockDim.x * gridDim.x; + + uint tmp = 0; + for (size_t i = idx; i < N; i += stride) { + for (size_t j = 0; j < ARRAY_SIZE; j++) { + tmp += src[i].data[j]; + } + } + + atomicAdd(dst, tmp); +} + +static bool hipPerfDevMemReadSpeed_test() { + d_uint16 *dSrc, *hSrc; + uint *dDst, *hDst; + hipStream_t stream; + ulong N = 4 * 1024 * 1024; + uint nBytes = N * sizeof(d_uint16); + + int deviceId = 0; + HIP_CHECK(hipSetDevice(deviceId)); + hipDeviceProp_t props = {0}; + HIP_CHECK(hipGetDeviceProperties(&props, deviceId)); + + INFO("info: running on bus " << "0x" << props.pciBusID << " " << + props.name << " with " << props.multiProcessorCount << " CUs \n"); + + const unsigned threadsPerBlock = 64; + const unsigned blocks = props.multiProcessorCount * 4; + + uint inputData = 0x1; + int nIter = 1000; + + hSrc = new d_uint16[nBytes]; + REQUIRE(hSrc != nullptr); + hDst = new uint; + REQUIRE(hDst != nullptr); + hDst[0] = 0; + + for (size_t i = 0; i < N; i++) { + for (int j = 0; j < ARRAY_SIZE; j++) { + hSrc[i].data[j] = inputData; + } + } + + HIP_CHECK(hipMalloc(&dSrc, nBytes)); + HIP_CHECK(hipMalloc(&dDst, sizeof(uint))); + + HIP_CHECK(hipStreamCreate(&stream)); + + HIP_CHECK(hipMemcpy(dSrc, hSrc, nBytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy(dDst, hDst, sizeof(uint), hipMemcpyHostToDevice)); + + hipLaunchKernelGGL(read_kernel, dim3(blocks), dim3(threadsPerBlock), + 0, stream, dSrc, N, dDst); + HIP_CHECK(hipGetLastError()); + HIP_CHECK(hipMemcpy(hDst, dDst, sizeof(uint), hipMemcpyDeviceToHost)); + HIP_CHECK(hipDeviceSynchronize()); + + if (hDst[0] != (nBytes / sizeof(uint))) { + INFO("hipPerfDevMemReadSpeed - Data validation failed for warm up run!" << + " expected " << nBytes / sizeof(uint) << " got " << hDst[0]); + return false; + } + + // measure performance based on host time + auto all_start = std::chrono::steady_clock::now(); + + for (int i = 0; i < nIter; i++) { + hipLaunchKernelGGL(read_kernel, dim3(blocks), dim3(threadsPerBlock), + 0, stream, dSrc, N, dDst); + HIP_CHECK(hipGetLastError()); + } + HIP_CHECK(hipDeviceSynchronize()); + + auto all_end = std::chrono::steady_clock::now(); + std::chrono::duration all_kernel_time = all_end - all_start; + + // read speed in GB/s + double perf = (static_cast(nBytes * nIter * (1e-09))) / + all_kernel_time.count(); + + INFO("hipPerfDevMemReadSpeed - info: average read speed of " << + perf << " GB/s " << "achieved for memory size of " << + nBytes / (1024 * 1024) << " MB"); + + delete [] hSrc; + delete hDst; + HIP_CHECK(hipFree(dSrc)); + HIP_CHECK(hipFree(dDst)); + HIP_CHECK(hipStreamDestroy(stream)); + return true; +} + +/** +* Test Description +* ------------------------ +*  - Verify hipPerfDevMemReadSpeed status. +* Test source +* ------------------------ +*  - perftests/memory/hipPerfDevMemReadSpeed.cc +* Test requirements +* ------------------------ +*  - HIP_VERSION >= 5.6 +*/ + +TEST_CASE("Perf_hipPerfDevMemReadSpeed_test") { + int numDevices = 0; + HIP_CHECK(hipGetDeviceCount(&numDevices)); + + if (numDevices <= 0) { + SUCCEED("Skipped testcase hipPerfDevMemReadSpeed as" + "there is no device to test."); + } else { + REQUIRE(true == hipPerfDevMemReadSpeed_test()); + } +} diff --git a/projects/hip-tests/catch/perftests/memory/hipPerfDevMemWriteSpeed.cc b/projects/hip-tests/catch/perftests/memory/hipPerfDevMemWriteSpeed.cc new file mode 100644 index 0000000000..ebc80d702b --- /dev/null +++ b/projects/hip-tests/catch/perftests/memory/hipPerfDevMemWriteSpeed.cc @@ -0,0 +1,147 @@ +/* +Copyright (c) 2023 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. +*/ + +/** +* @addtogroup hipMemcpyKernel hipMemcpyKernel +* @{ +* @ingroup perfMemoryTest +* `hipMemcpy(void* dst, const void* src, size_t count, hipMemcpyKind kind)` - +* Copies data between host and device. +*/ + +#include + +#define ARRAY_SIZE 16 + +typedef struct d_uint16 { + uint data[ARRAY_SIZE]; +} d_uint16; + +__global__ void write_kernel(d_uint16 *dst, ulong N, d_uint16 pval) { + size_t idx = (blockIdx.x * blockDim.x + threadIdx.x); + size_t stride = blockDim.x * gridDim.x; + for (size_t i = idx; i < N; i += stride) { + dst[i] = pval; + } +} + +static bool hipPerfDevMemWriteSpeed_test() { + d_uint16 *dDst, *hDst; + ulong N = 4 * 1024 * 1024; + uint nBytes = N * sizeof(d_uint16); + + uint inputData = 0xabababab; + int nIter = 1000; + d_uint16 pval; + + int deviceId = 0; + HIP_CHECK(hipSetDevice(deviceId)); + hipDeviceProp_t props = {0}; + HIP_CHECK(hipGetDeviceProperties(&props, deviceId)); + + INFO("info: running on bus " << "0x" << props.pciBusID << " " << + props.name << " with " << props.multiProcessorCount << " CUs \n"); + + const unsigned threadsPerBlock = 64; + const unsigned blocks = props.multiProcessorCount * 4; + + for (int i = 0; i < ARRAY_SIZE; i++) { + pval.data[i] = inputData; + } + + hDst = new d_uint16[nBytes]; + REQUIRE(hDst != nullptr); + + for (size_t i = 0; i < N; i++) { + for (size_t j = 0; j < ARRAY_SIZE; j++) { + hDst[i].data[j] = 0; + } + } + + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + HIP_CHECK(hipMalloc(&dDst, nBytes)); + hipLaunchKernelGGL(write_kernel, dim3(blocks), dim3(threadsPerBlock), + 0, stream, dDst, N, pval); + HIP_CHECK(hipGetLastError()); + HIP_CHECK(hipMemcpy(hDst, dDst, nBytes , hipMemcpyDeviceToHost)); + HIP_CHECK(hipDeviceSynchronize()); + + for (uint i = 0; i < N; i++) { + for (uint j = 0; j < ARRAY_SIZE; j++) { + if (hDst[i].data[j] != inputData) { + INFO("hipPerfDevMemWriteSpeed - Data validation failed for warm up run!" + << "at index i: " << i << " element j: " << j << + "expected " << inputData << " but got " << hDst[i].data[j]); + return false; + } + } + } + + // measure performance based on host time + auto all_start = std::chrono::steady_clock::now(); + + for (int i = 0; i < nIter; i++) { + hipLaunchKernelGGL(write_kernel, dim3(blocks), dim3(threadsPerBlock), + 0, stream, dDst, N, pval); + HIP_CHECK(hipGetLastError()); + } + HIP_CHECK(hipDeviceSynchronize()); + + auto all_end = std::chrono::steady_clock::now(); + std::chrono::duration all_kernel_time = all_end - all_start; + + // read speed in GB/s + double perf = (static_cast(nBytes * nIter * (1e-09))) / + all_kernel_time.count(); + + INFO("hipPerfDevMemReadSpeed - info: average write speed of " << + perf << " GB/s " << "achieved for memory size of " << + nBytes / (1024 * 1024) << " MB"); + + delete [] hDst; + HIP_CHECK(hipFree(dDst)); + HIP_CHECK(hipStreamDestroy(stream)); + return true; +} + +/** +* Test Description +* ------------------------ +*  - Verify hipPerfDevMemWriteSpeed status. +* Test source +* ------------------------ +*  - perftests/memory/hipPerfDevMemWriteSpeed.cc +* Test requirements +* ------------------------ +*  - HIP_VERSION >= 5.6 +*/ + +TEST_CASE("Perf_hipPerfDevMemWriteSpeed_test") { + int numDevices = 0; + HIP_CHECK(hipGetDeviceCount(&numDevices)); + + if (numDevices <= 0) { + SUCCEED("Skipped testcase hipPerfDevMemWriteSpeed as" + "there is no device to test."); + } else { + REQUIRE(true == hipPerfDevMemWriteSpeed_test()); + } +} diff --git a/projects/hip-tests/catch/perftests/memory/hipPerfHostNumaAlloc.cc b/projects/hip-tests/catch/perftests/memory/hipPerfHostNumaAlloc.cc new file mode 100644 index 0000000000..4c7e4b7a86 --- /dev/null +++ b/projects/hip-tests/catch/perftests/memory/hipPerfHostNumaAlloc.cc @@ -0,0 +1,191 @@ +/* +Copyright (c) 2023 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. +*/ + +/** +* @addtogroup hipMemcpyKernel hipMemcpyKernel +* @{ +* @ingroup perfMemoryTest +* `hipMemcpy(void* dst, const void* src, size_t count, hipMemcpyKind kind)` - +* Copies data between host and device. +*/ + +#include +#include + +// To run it correctly, we must not export HIP_VISIBLE_DEVICES. +// And we must explicitly link libnuma because of numa api move_pages(). +#define NUM_PAGES 4 +char *h = nullptr; +char *d_h = nullptr; +char *m = nullptr; +char *d_m = nullptr; +int page_size = 1024; + +const int mode[] = { MPOL_DEFAULT, MPOL_BIND, MPOL_PREFERRED, MPOL_INTERLEAVE }; +const char* modeStr[] = { "MPOL_DEFAULT", "MPOL_BIND", + "MPOL_PREFERRED", "MPOL_INTERLEAVE" }; + +std::string exeCommand(const char* cmd) { + std::array buff; + std::string result; + std::unique_ptr pipe(popen(cmd, "r"), pclose); + if (!pipe) { + return result; + } + while (fgets(buff.data(), buff.size(), pipe.get()) != nullptr) { + result += buff.data(); + } + return result; +} + +int getCpuAgentCount() { + const char* cmd = + "cat /proc/cpuinfo | grep \"physical id\" | sort | uniq | wc -l"; + int cpuAgentCount = std::atoi(exeCommand(cmd).c_str()); + return cpuAgentCount; +} + +bool test(int cpuId, int gpuId, int numaMode, unsigned int hostMallocflags) { + void *pages[NUM_PAGES]; + int status[NUM_PAGES]; + int nodes[NUM_PAGES]; + int ret_code; + + INFO("set cpu " << cpuId << ", gpu " << gpuId << ", numaMode " + << numaMode << ", hostMallocflags " << hostMallocflags << "\n"); + + if (cpuId >= 0) { + unsigned long nodeMask = 1 << cpuId; //NOLINT + unsigned long maxNode = sizeof(nodeMask) * 8; //NOLINT + if (set_mempolicy(numaMode, numaMode == MPOL_DEFAULT ? NULL : &nodeMask, + numaMode == MPOL_DEFAULT ? 0 : maxNode) == -1) { + WARN("set_mempolicy() failed with err " << errno << "\n"); + return false; + } + } + + if (gpuId >= 0) { + HIP_CHECK(hipSetDevice(gpuId)); + } + + posix_memalign(reinterpret_cast(&m), page_size, page_size*NUM_PAGES); + HIP_CHECK(hipHostRegister(m, page_size * NUM_PAGES, hipHostRegisterMapped)); + HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&d_m), m, 0)); + + status[0] = -1; + pages[0] = m; + for (int i = 1; i < NUM_PAGES; i++) { + pages[i] = reinterpret_cast(pages[0]) + page_size; + } + + ret_code = move_pages(0, NUM_PAGES, pages, NULL, status, 0); + INFO("Memory (malloc) ret " << ret_code << " at " << m << + " (dev " << d_m << "%p) is at node: "); + for (int i = 0; i < NUM_PAGES; i++) { + INFO(status[i]); // Don't verify as it's out of our control + } + INFO("\n"); + + HIP_CHECK(hipHostMalloc(reinterpret_cast(&h), + page_size*NUM_PAGES, hostMallocflags)); + pages[0] = h; + for (int i = 1; i < NUM_PAGES; i++) { + pages[i] = reinterpret_cast(pages[0]) + page_size; + } + ret_code = move_pages(0, NUM_PAGES, pages, NULL, status, 0); + d_h = nullptr; + if (hostMallocflags & hipHostMallocMapped) { + HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&d_h), h, 0)); + INFO("Memory (hipHostMalloc) ret " << ret_code << " at " << h + << " (dev " << d_h << ") is at node: "); + } else { + INFO("Memory (hipHostMalloc) ret " << ret_code << " at " + << h << " is at node: "); + } + for (int i = 0; i < NUM_PAGES; i++) { + INFO(status[i]); // Always print it even if it's wrong. Verify later + } + INFO("\n"); + + HIP_CHECK(hipHostFree(reinterpret_cast(h))); + HIP_CHECK(hipHostUnregister(m)); + free(m); + + if (cpuId >= 0 && (numaMode == MPOL_BIND || numaMode == MPOL_PREFERRED)) { + for (int i = 0; i < NUM_PAGES; i++) { + if (status[i] != cpuId) { // Now verify + WARN("Failed at " << i << " status[i] = " << status[i] + << " cpuId " << cpuId << "\n"); + return false; + } + } + } + return true; +} + +bool runTest(const int &cpuCount, const int &gpuCount, + unsigned int hostMallocflags, const std::string &str) { + INFO("Test- " << str.c_str() << "\n"); + + for (int m = 0; m < sizeof(mode) / sizeof(mode[0]); m++) { + INFO("Testing " << modeStr[m] << "\n"); + + for (int i = 0; i < cpuCount; i++) { + for (int j = 0; j < gpuCount; j++) { + if (!test(i, j, mode[m], hostMallocflags)) { + return false; + } + } + } + } + return true; +} + +/** +* Test Description +* ------------------------ +*  - Verify hipPerfHostNumaAlloc status. +* Test source +* ------------------------ +*  - perftests/memory/hipPerfHostNumaAlloc.cc +* Test requirements +* ------------------------ +*  - HIP_VERSION >= 5.6 +*/ + +TEST_CASE("Perf_hipPerfHostNumaAlloc_test") { + int gpuCount = 0; + HIP_CHECK(hipGetDeviceCount(&gpuCount)); + int cpuCount = getCpuAgentCount(); + INFO("Cpu count " << cpuCount << ", Gpu count " << gpuCount << "\n"); + + if (cpuCount < 0 || gpuCount < 0) { + SUCCEED("Skipped testcase hipPerfHostNumaAlloc as " + "there is no device to test.\n"); + return; + } + + REQUIRE(true == runTest(cpuCount, gpuCount, + hipHostMallocDefault | hipHostMallocNumaUser, + "Testing hipHostMallocDefault | hipHostMallocNumaUser......")); + + REQUIRE(true == runTest(cpuCount, gpuCount, + hipHostMallocMapped | hipHostMallocNumaUser, + "Testing hipHostMallocMapped | hipHostMallocNumaUser.......")); +} diff --git a/projects/hip-tests/catch/perftests/memory/hipPerfMemFill.cc b/projects/hip-tests/catch/perftests/memory/hipPerfMemFill.cc new file mode 100644 index 0000000000..ae10829d17 --- /dev/null +++ b/projects/hip-tests/catch/perftests/memory/hipPerfMemFill.cc @@ -0,0 +1,545 @@ +/* + Copyright (c) 2023 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. + */ + +/** +* @addtogroup hipMemcpyKernel hipMemcpyKernel +* @{ +* @ingroup perfMemoryTest +* `hipMemcpy(void* dst, const void* src, size_t count, hipMemcpyKind kind)` - +* Copies data between host and device. +*/ + +#include + +#define SIMPLY_ASSIGN 0 +#define USE_HIPTEST_SETNUMBLOCKS 0 + +template +__global__ void vec_fill(T *x, T coef, int N) { + const int istart = threadIdx.x + blockIdx.x * blockDim.x; + const int ishift = blockDim.x * gridDim.x; + for (int i = istart; i < N; i += ishift) { +#if SIMPLY_ASSIGN + x[i] = coef; +#else + x[i] = coef * i; +#endif + } +} + +__device__ void print_log(int i, double value, double expected) { + printf("failed at %d: val=%g, expected=%g\n", i, value, expected); +} + +__device__ void print_log(int i, int value, int expected) { + printf("failed at %d: val=%d, expected=%d\n", i, value, expected); +} + +template +__global__ void vec_verify(T *x, T coef, int N) { + const int istart = threadIdx.x + blockIdx.x * blockDim.x; + const int ishift = blockDim.x * gridDim.x; + for (int i = istart; i < N; i += ishift) { +#if SIMPLY_ASSIGN + if (x[i] != coef) { + print_log(i, x[i], coef); + } +#else + if (x[i] != coef * i) { + print_log(i, x[i], coef * i); + } +#endif + } +} + +template +__global__ void daxpy(T *__restrict__ x, T *__restrict__ y, + const T coef, int Niter, int N) { + const int istart = threadIdx.x + blockIdx.x * blockDim.x; + const int ishift = blockDim.x * gridDim.x; + for (int iter = 0; iter < Niter; ++iter) { + T iv = coef * iter; + for (int i = istart; i < N; i += ishift) + y[i] = iv * x[i] + y[i]; + } +} + +template +class hipPerfMemFill { + private: + static constexpr int NUM_START = 27; + static constexpr int NUM_SIZE = 5; + static constexpr int NUM_ITER = 10; + static constexpr double NUM_1GB = 1024.0 * 1024.0 * 1024.0; + size_t totalSizes_[NUM_SIZE]; + hipDeviceProp_t props_; + const T coef_ = getCoefficient(3.14159); + const unsigned int threadsPerBlock_ = 64; + unsigned int blocksPerCU_; + + public: + hipPerfMemFill() { + for (int i = 0; i < NUM_SIZE; i++) { + // 128M, 256M, 512M, 1024M, 2048M + totalSizes_[i] = 1ull << (i + NUM_START); + } + } + + ~hipPerfMemFill() { } + + bool supportLargeBar() { + return props_.isLargeBar != 0; + } + + bool supportManagedMemory() { + return props_.managedMemory != 0; + } + + const T getCoefficient(double val) { + return static_cast(val); + } + + void setHostBuffer(T *A, T val, size_t size) { + size_t len = size / sizeof(T); + for (int i = 0; i < len; i++) { + A[i] = val; + } + } + + bool open(int deviceId) { + int nGpu = 0; + HIP_CHECK(hipGetDeviceCount(&nGpu)); + if (nGpu < 1) { + printf("No GPU!"); + return false; + } else if (deviceId >= nGpu) { + printf("Info: wrong GPU Id %d\n", deviceId); + return false; + } + + HIP_CHECK(hipSetDevice(deviceId)); + memset(&props_, 0, sizeof(props_)); + HIP_CHECK(hipGetDeviceProperties(&props_, deviceId)); + blocksPerCU_ = props_.multiProcessorCount * 4; + + std::cout << "Info: running on device: id: " << deviceId << ", bus: 0x" + << props_.pciBusID << " " << props_.name << " with " + << props_.multiProcessorCount << " CUs, large bar: " + << supportLargeBar() << ", managed memory: " << supportManagedMemory() + << ", DeviceMallocFinegrained: " << supportDeviceMallocFinegrained() + << std::endl; + return true; + } + + void log_host(const char* title, double GBytes, double sec) { + std::cout << title << " [" << std::setw(7) << GBytes << " GB]: cost " + << std::setw(10) << sec << " s in bandwidth " << std::setw(10) + << GBytes / sec << " [GB/s]" << std::endl; + } + + void log_kernel(const char* title, double GBytes, double sec, + double sec_hv, double sec_kv) { + std::cout << title << " [" << std::setw(7) << GBytes << " GB]: cost " + << std::setw(10) << sec << " s in bandwidth " << std::setw(10) + << GBytes / sec << " [GB/s]" << ", hostVerify cost " + << std::setw(10) << sec_hv << " s in bandwidth " << std::setw(10) + << GBytes / sec_hv << " [GB/s]" << ", kernelVerify cost " + << std::setw(10) << sec_kv << " s in bandwidth " << std::setw(10) + << GBytes / sec_kv << " [GB/s]" << std::endl; + } + + void hostFill(size_t size, T *data, T coef, double *sec) { + size_t num = size / sizeof(T); // Size of elements + auto start = std::chrono::steady_clock::now(); + for (int i = 0; i < num; ++i) { +#if SIMPLY_ASSIGN + data[i] = coef; +#else + data[i] = coef * i; +#endif + } + auto end = std::chrono::steady_clock::now(); + std::chrono::duration diff = end - start; // in second + *sec = diff.count(); + } + + void kernelFill(size_t size, T *data, T coef, double *sec) { + size_t num = size / sizeof(T); // Size of elements + unsigned blocks = setNumBlocks(num); + + // kernel will be loaded first time + hipLaunchKernelGGL(HIP_KERNEL_NAME(vec_fill), dim3(blocks), + dim3(threadsPerBlock_), 0, 0, data, 0, num); + HIP_CHECK(hipDeviceSynchronize()); + + auto start = std::chrono::steady_clock::now(); + + for (int iter = 0; iter < NUM_ITER; ++iter) { + hipLaunchKernelGGL(HIP_KERNEL_NAME(vec_fill), dim3(blocks), + dim3(threadsPerBlock_), 0, 0, data, coef, num); + } + HIP_CHECK(hipDeviceSynchronize()); + + auto end = std::chrono::steady_clock::now(); + std::chrono::duration diff = end - start; // in second + *sec = diff.count() / NUM_ITER; // in second + } + + void hostVerify(size_t size, T *data, T coef, double *sec) { + size_t num = size / sizeof(T); // Size of elements + auto start = std::chrono::steady_clock::now(); + for (int i = 0; i < num; ++i) { +#if SIMPLY_ASSIGN + if (data[i] != coef) { + std::cout << "hostVerify failed: i=" << i << ", data[i]=" << data[i] + << ", expected=" << coef << std::endl; + REQUIRE(false); + } +#else + if (data[i] != coef * i) { + std::cout << "hostVerify failed: i=" << i << ", data[i]=" << data[i] + << ", expected=" << coef * i << std::endl; + REQUIRE(false); + } +#endif + } + auto end = std::chrono::steady_clock::now(); + std::chrono::duration diff = end - start; // in second + *sec = diff.count(); + } + + void kernelVerify(size_t size, T *data, T coef, double *sec) { + size_t num = size / sizeof(T); // Size of elements + unsigned blocks = setNumBlocks(num); + + // kernel will be loaded first time + hipLaunchKernelGGL(HIP_KERNEL_NAME(vec_verify), dim3(blocks), + dim3(threadsPerBlock_), 0, 0, data, coef, num); + HIP_CHECK(hipDeviceSynchronize()); + + // Now all data verified. The following is to test bandwidth. + auto start = std::chrono::steady_clock::now(); + + for (int iter = 0; iter < NUM_ITER; ++iter) { + hipLaunchKernelGGL(HIP_KERNEL_NAME(vec_verify), dim3(blocks), + dim3(threadsPerBlock_), 0, 0, data, coef, num); + } + HIP_CHECK(hipDeviceSynchronize()); + + auto end = std::chrono::steady_clock::now(); + std::chrono::duration diff = end - start; // in second + *sec = diff.count() / NUM_ITER; // in second + } + + bool testLargeBarDeviceMemoryHostFill(size_t size) { + if (!supportLargeBar()) { + return false; + } + + double GBytes = static_cast(size) / NUM_1GB; + + T *A; + HIP_CHECK(hipMalloc(&A, size)); + double sec = 0; + hostFill(size, A, coef_, &sec); // Cpu can access device mem in LB + HIP_CHECK(hipFree(A)); + + log_host("Largebar: host fill", GBytes, sec); + return true; + } + + bool testLargeBar() { + if (!supportLargeBar()) { + return false; + } + + std::cout << "Test large bar device memory host filling" << std::endl; + for (int i = 0; i < NUM_SIZE; i++) { + if (!testLargeBarDeviceMemoryHostFill(totalSizes_[i])) { + return false; + } + } + + return true; + } + + bool testManagedMemoryHostFill(size_t size) { + if (!supportManagedMemory()) { + return false; + } + double GBytes = static_cast(size) / NUM_1GB; + + T *A; + HIP_CHECK(hipMallocManaged(&A, size)); + double sec = 0; + hostFill(size, A, coef_, &sec); // Cpu can access HMM mem + HIP_CHECK(hipFree(A)); + + log_host("Managed: host fill", GBytes, sec); + return true; + } + + bool testManagedMemoryKernelFill(size_t size) { + if (!supportManagedMemory()) { + return false; + } + double GBytes = static_cast(size) / NUM_1GB; + + T *A; + HIP_CHECK(hipMallocManaged(&A, size)); + + double sec = 0, sec_hv = 0, sec_kv = 0; + kernelFill(size, A, coef_, &sec); + // Managed memory can be verified by host + hostVerify(size, A, coef_, &sec_hv); + kernelVerify(size, A, coef_, &sec_kv); + HIP_CHECK(hipFree(A)); + + log_kernel("Managed: kernel fill", GBytes, sec, sec_hv, sec_kv); + + return true; + } + + bool testManagedMemory() { + if (!supportManagedMemory()) { + return false; + } + + std::cout << "Test managed memory host filling" << std::endl; + for (int i = 0; i < NUM_SIZE; i++) { + if (!testManagedMemoryHostFill(totalSizes_[i])) { + return false; + } + } + + std::cout << "Test managed memory kernel filling" << std::endl; + for (int i = 0; i < NUM_SIZE; i++) { + if (!testManagedMemoryKernelFill(totalSizes_[i])) { + return false; + } + } + + return true; + } + + bool testHostMemoryHostFill(size_t size, unsigned int flags) { + double GBytes = static_cast(size) / NUM_1GB; + T *A; + HIP_CHECK(hipHostMalloc(&A, size, flags)); + double sec = 0; + hostFill(size, A, coef_, &sec); + HIP_CHECK(hipHostFree(A)); + + log_host("Host: host fill", GBytes, sec); + return true; + } + + bool testHostMemoryKernelFill(size_t size, unsigned int flags) { + double GBytes = static_cast(size) / NUM_1GB; + + T *A; + HIP_CHECK(hipHostMalloc(reinterpret_cast(&A), size, flags)); + double sec = 0, sec_hv = 0, sec_kv = 0; + kernelFill(size, A, coef_, &sec); + hostVerify(size, A, coef_, &sec_hv); + kernelVerify(size, A, coef_, &sec_kv); + HIP_CHECK(hipHostFree(A)); + + log_kernel("Host: kernel fill", GBytes, sec, sec_hv, sec_kv); + return true; + } + + bool testHostMemory() { + std::cout << "Test coherent host memory host filling" << std::endl; + for (int i = 0; i < NUM_SIZE; i++) { + if (!testHostMemoryHostFill(totalSizes_[i], hipHostMallocCoherent)) { + return false; + } + } + + std::cout << "Test non-coherent host memory host filling" << std::endl; + for (int i = 0; i < NUM_SIZE; i++) { + if (!testHostMemoryHostFill(totalSizes_[i], hipHostMallocNonCoherent)) { + return false; + } + } + + std::cout << "Test coherent host memory kernel filling" << std::endl; + for (int i = 0; i < NUM_SIZE; i++) { + if (!testHostMemoryKernelFill(totalSizes_[i], hipHostMallocCoherent)) { + return false; + } + } + + std::cout << "Test non-coherent host memory kernel filling" << std::endl; + for (int i = 0; i < NUM_SIZE; i++) { + if (!testHostMemoryKernelFill(totalSizes_[i], hipHostMallocNonCoherent)) { + return false; + } + } + + return true; + } + + /* This function should be via device attribute query*/ + bool supportDeviceMallocFinegrained() { +#ifdef __HIP_PLATFORM_AMD__ + T *A = nullptr; + hipError_t err; + err = hipExtMallocWithFlags(reinterpret_cast(&A), sizeof(T), + hipDeviceMallocFinegrained); + if (err || !A) { + return false; + } + HIP_CHECK(hipFree(A)); + return true; +#else + return false; +#endif + } + + unsigned int setNumBlocks(size_t size) { + size_t num = size/sizeof(T); + +#if USE_HIPTEST_SETNUMBLOCKS + return HipTest::setNumBlocks(blocksPerCU_, threadsPerBlock_, num); +#else + return (num + threadsPerBlock_ - 1) / threadsPerBlock_; +#endif + } + +#ifdef __HIP_PLATFORM_AMD__ + bool testExtDeviceMemoryHostFill(size_t size, unsigned int flags) { + double GBytes = static_cast(size) / NUM_1GB; + + T *A = nullptr; + HIP_CHECK(hipExtMallocWithFlags(reinterpret_cast(&A), + size, flags)); + if (!A) { + std::cout << "failed hipExtMallocWithFlags() with size =" << + size << " flags="<< std::hex << flags << std::endl; + return false; + } + + double sec = 0; + hostFill(size, A, coef_, &sec); // Cpu can access this mem + HIP_CHECK(hipFree(A)); + + log_host("ExtDevice: host fill", GBytes, sec); + return true; + } + + bool testExtDeviceMemoryKernelFill(size_t size, unsigned int flags) { + double GBytes = static_cast(size) / NUM_1GB; + + T *A = nullptr; + HIP_CHECK(hipExtMallocWithFlags(reinterpret_cast(&A), + size, flags)); + if (!A) { + std::cout << "failed hipExtMallocWithFlags() with size =" << + size << " flags=" << std::hex << flags << std::endl; + return false; + } + + double sec = 0, sec_hv = 0, sec_kv = 0; + kernelFill(size, A, coef_, &sec); + // Fine grained device memory can be verified by host + hostVerify(size, A, coef_, &sec_hv); + kernelVerify(size, A, coef_, &sec_kv); + HIP_CHECK(hipFree(A)); + + log_kernel("ExtDevice: kernel fill", GBytes, sec, sec_hv, sec_kv); + + return true; + } + + bool testExtDeviceMemory() { + std::cout << "Test fine grained device memory host filling" + << std::endl; + for (int i = 0; i < NUM_SIZE; i++) { + if (!testExtDeviceMemoryHostFill(totalSizes_[i], + hipDeviceMallocFinegrained)) { + return false; + } + } + + std::cout << "Test fine grained device memory kernel filling" + << std::endl; + for (int i = 0; i < NUM_SIZE; i++) { + if (!testExtDeviceMemoryKernelFill(totalSizes_[i], + hipDeviceMallocFinegrained)) { + return false; + } + } + + return true; + } +#endif + + bool run() { + if (supportLargeBar()) { + if (!testLargeBar()) { + return false; + } + } + + if (supportManagedMemory()) { + if (!testManagedMemory()) { + return false; + } + } + + if (!testHostMemory()) { + return false; + } + +#ifdef __HIP_PLATFORM_AMD__ + if (supportDeviceMallocFinegrained()) { + if (!testExtDeviceMemory()) { + return false; + } + } +#endif + return true; + } +}; + +/** +* Test Description +* ------------------------ +*  - Verify hipPerfMemFill status. +* Test source +* ------------------------ +*  - perftests/memory/hipPerfMemFill.cc +* Test requirements +* ------------------------ +*  - HIP_VERSION >= 5.6 +*/ + +TEST_CASE("Perf_hipPerfMemFill_test") { + std::cout << "Test int" << std::endl; + hipPerfMemFill hipPerfMemFillInt; + REQUIRE(true == hipPerfMemFillInt.open(0)); + REQUIRE(true == hipPerfMemFillInt.run()); + + std::cout << "Test double" << std::endl; + hipPerfMemFill hipPerfMemFillDouble; + REQUIRE(true == hipPerfMemFillDouble.open(0)); + REQUIRE(true == hipPerfMemFillDouble.run()); +} diff --git a/projects/hip-tests/catch/perftests/memory/hipPerfMemMallocCpyFree.cc b/projects/hip-tests/catch/perftests/memory/hipPerfMemMallocCpyFree.cc new file mode 100644 index 0000000000..74d9adc346 --- /dev/null +++ b/projects/hip-tests/catch/perftests/memory/hipPerfMemMallocCpyFree.cc @@ -0,0 +1,144 @@ +/* +Copyright (c) 2023 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. +*/ + +/** +* @addtogroup hipMemcpy hipMemcpy +* @{ +* @ingroup perfMemoryTest +* `hipMemcpy(void* dst, const void* src, size_t count, hipMemcpyKind kind)` - +* Copies data between host and device. +*/ + +#include +#include + +#define NUM_SIZE 19 // size up to 16M +#define NUM_ITER 500 // Total GPU memory up to 16M*500=8G + +void valSet(int* A, int val, size_t size) { + size_t len = size / sizeof(int); + for (int i = 0; i < len; i++) { + A[i] = val; + } +} + +void setup(size_t *size, int *num, int **pA, const size_t totalGlobalMem) { + for (int i = 0; i < *num; i++) { + size[i] = 1 << (i + 6); + if ((NUM_ITER + 1) * size[i] > totalGlobalMem) { + *num = i; + break; + } + } + *pA = reinterpret_cast(malloc(size[*num - 1])); + valSet(*pA, 1, size[*num - 1]); +} + +void testInit(size_t size, int *A) { + int *Ad; + + clock_t start = clock(); + HIP_CHECK(hipMalloc(&Ad, size)); // hip::init() will be called + clock_t end = clock(); + double uS = (end - start) * 1000000. / CLOCKS_PER_SEC; + INFO("Initial: hipMalloc(" << size << ") cost " << uS << "us" << "\n"); + + start = clock(); + HIP_CHECK(hipMemcpy(Ad, A, size, hipMemcpyHostToDevice)); + HIP_CHECK(hipDeviceSynchronize()); + end = clock(); + uS = (end - start) * 1000000. / CLOCKS_PER_SEC; + INFO("hipMemcpy(" << size << ") cost " << uS << "us" << "\n"); + + start = clock(); + HIP_CHECK(hipFree(Ad)); + end = clock(); + uS = (end - start) * 1000000. / CLOCKS_PER_SEC; + INFO("hipFree(" << size << ") cost " << uS << "us" << "\n"); +} + +static bool hipPerfMemMallocCpyFree_test() { + double uS; + clock_t start, end; + size_t size[NUM_SIZE] = { 0 }; + int *Ad[NUM_ITER] = { nullptr }; + int *A; + hipDeviceProp_t props; + memset(&props, 0, sizeof(props)); + HIP_CHECK(hipGetDeviceProperties(&props, 0)); + INFO("totalGlobalMem: " << props.totalGlobalMem << "\n"); + + int num = NUM_SIZE; + setup(size, &num, &A, props.totalGlobalMem); + testInit(size[0], A); + + for (int i = 0; i < num; i++) { + start = clock(); + for (int j = 0; j < NUM_ITER; j++) { + HIP_CHECK(hipMalloc(&Ad[j], size[i])); + } + end = clock(); + uS = (end - start) * 1000000. / (NUM_ITER * CLOCKS_PER_SEC); + INFO("hipMalloc(" << size[i] << ") cost " << uS << "us" << "\n"); + + start = clock(); + for (int j = 0; j < NUM_ITER; j++) { + HIP_CHECK(hipMemcpy(Ad[j], A, size[i], hipMemcpyHostToDevice)); + } + HIP_CHECK(hipDeviceSynchronize()); + end = clock(); + uS = (end - start) * 1000000. / (NUM_ITER * CLOCKS_PER_SEC); + INFO("hipMemcpy(" << size[i] << ") cost " << uS << "us" << "\n"); + + start = clock(); + for (int j = 0; j < NUM_ITER; j++) { + HIP_CHECK(hipFree(Ad[j])); + Ad[j] = nullptr; + } + end = clock(); + double uS = (end - start) * 1000000. / (NUM_ITER * CLOCKS_PER_SEC); + INFO("hipFree(" << size[i] << ") cost " << uS << "us" << "\n"); + } + free(A); + return true; +} + +/** +* Test Description +* ------------------------ +*  - Verify hipPerfMemMallocCpyFree status. +* Test source +* ------------------------ +*  - perftests/memory/hipPerfMemMallocCpyFree.cc +* Test requirements +* ------------------------ +*  - HIP_VERSION >= 5.6 +*/ + +TEST_CASE("Perf_hipPerfMemMallocCpyFree_test") { + int numDevices = 0; + HIP_CHECK(hipGetDeviceCount(&numDevices)); + + if (numDevices <= 0) { + SUCCEED("Skipped testcase hipPerfDevMemReadSpeed as" + "there is no device to test."); + } else { + REQUIRE(true == hipPerfMemMallocCpyFree_test()); + } +} diff --git a/projects/hip-tests/catch/perftests/memory/hipPerfMemcpy.cc b/projects/hip-tests/catch/perftests/memory/hipPerfMemcpy.cc new file mode 100644 index 0000000000..bc40760411 --- /dev/null +++ b/projects/hip-tests/catch/perftests/memory/hipPerfMemcpy.cc @@ -0,0 +1,118 @@ +/* + Copyright (c) 2023 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. + */ + +/** +* @addtogroup hipMemcpy hipMemcpy +* @{ +* @ingroup perfMemoryTest +* `hipMemcpy(void* dst, const void* src, size_t count, hipMemcpyKind kind)` - +* Copies data between host and device. +*/ + +#include + +#define NUM_SIZE 8 +#define NUM_ITER 0x40000 + +class hipPerfMemcpy { + private: + unsigned int numBuffers_; + size_t totalSizes_[NUM_SIZE]; + void setHostBuffer(int *A, int val, size_t size); + public: + hipPerfMemcpy(); + ~hipPerfMemcpy() {} + bool run(unsigned int numTests); +}; + +hipPerfMemcpy::hipPerfMemcpy() : numBuffers_(0) { + for (int i = 0; i < NUM_SIZE; i++) { + totalSizes_[i] = 1 << (i + 6); + } +} + +void hipPerfMemcpy::setHostBuffer(int *A, int val, size_t size) { + size_t len = size / sizeof(int); + for (int i = 0; i < len; i++) { + A[i] = val; + } +} + +bool hipPerfMemcpy::run(unsigned int numTests) { + int *A, *Ad; + A = new int[totalSizes_[numTests]]; + setHostBuffer(A, 1, totalSizes_[numTests]); + HIP_CHECK(hipMalloc(&Ad, totalSizes_[numTests])); + + // measure performance based on host time + auto all_start = std::chrono::steady_clock::now(); + + for (int j = 0; j < NUM_ITER; j++) { + HIP_CHECK(hipMemcpy(Ad, A, totalSizes_[numTests], hipMemcpyHostToDevice)); + } + + HIP_CHECK(hipDeviceSynchronize()); + + auto all_end = std::chrono::steady_clock::now(); + std::chrono::duration diff = all_end - all_start; + + INFO("hipPerfMemcpy[" << numTests << "] " << "Host to Device copy took " + << diff.count() / NUM_ITER << " sec for memory size of " << + totalSizes_[numTests] << " Bytes."); + + delete [] A; + HIP_CHECK(hipFree(Ad)); + + return true; +} + +/** +* Test Description +* ------------------------ +*  - Verify hipPerfMemcpy status. +* Test source +* ------------------------ +*  - perftests/memory/hipPerfMemcpy.cc +* Test requirements +* ------------------------ +*  - HIP_VERSION >= 5.6 +*/ + +TEST_CASE("Perf_hipPerfMemcpy_test") { + int numDevices = 0; + HIP_CHECK(hipGetDeviceCount(&numDevices)); + + if (numDevices <= 0) { + SUCCEED("Skipped testcase hipPerfMemcpy as there is no device to test."); + } else { + int deviceId = 0; + HIP_CHECK(hipSetDevice(deviceId)); + hipDeviceProp_t props = {0}; + HIP_CHECK(hipGetDeviceProperties(&props, deviceId)); + + INFO("info: running on bus " << "0x" << props.pciBusID << " " << + props.name << " with " << props.multiProcessorCount << " CUs " + << " and device id: " << deviceId); + + hipPerfMemcpy hipPerfMemcpy; + for (auto testCase = 0; testCase < NUM_SIZE; testCase++) { + REQUIRE(true == hipPerfMemcpy.run(testCase)); + } + } +} diff --git a/projects/hip-tests/catch/perftests/memory/hipPerfMemset.cc b/projects/hip-tests/catch/perftests/memory/hipPerfMemset.cc new file mode 100644 index 0000000000..7ae952e02c --- /dev/null +++ b/projects/hip-tests/catch/perftests/memory/hipPerfMemset.cc @@ -0,0 +1,419 @@ +/* + Copyright (c) 2023 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. + */ + +/** +* @addtogroup hipMemsetKernel hipMemsetKernel +* @{ +* @ingroup perfMemoryTest +* `hipMemset(void* devPtr, int value, size_t count)` - +* Initializes or sets device memory to a value. +*/ + +#include + +static size_t typeSizeList[] = { + 1, 2, 4, 8, 16, 32, 64, 128, +}; + +static unsigned int sizeList[] = { + 256, 512, 1024, 2048, 4096, 8192, +}; + +static unsigned int eleNumList[] = { + 0x0020000, 0x0080000, 0x0200000, 0x0800000, 0x2000000, +}; + +typedef struct _dataType { + char memsetval = 0x42; + char memsetD8val = 0xDE; + int16_t memsetD16val = 0xDEAD; + int memsetD32val = 0xDEADBEEF; +}dataType; + +#define NUM_ITER 100 + +enum MemsetType { + hipMemsetTypeDefault, + hipMemsetTypeD8, + hipMemsetTypeD16, + hipMemsetTypeD32 +}; + +class hipPerfMemset { + private: + unsigned int bufSize_; + unsigned int num_typeSize_; + unsigned int num_elements_; + size_t testTypeSize_; + unsigned int testNumEle_; + unsigned int _numSubTests = 0; + unsigned int _numSubTests2D = 0; + unsigned int _numSubTests3D = 0; + unsigned int num_sizes_ = 0; + + public: + hipPerfMemset() { + num_typeSize_ = sizeof(typeSizeList) / sizeof(size_t); + num_elements_ = sizeof(eleNumList) / sizeof(unsigned int); + _numSubTests = num_elements_ * num_typeSize_; + + num_sizes_ = sizeof(sizeList) / sizeof(unsigned int); + _numSubTests2D = num_sizes_; + _numSubTests3D = _numSubTests2D; + } + + ~hipPerfMemset() {} + + bool open(int deviceID); + + template + void run1D(unsigned int test, T memsetval, enum MemsetType type, bool async); + + template + void run2D(unsigned int test, T memsetval, enum MemsetType type, bool async); + + template + void run3D(unsigned int test, T memsetval, enum MemsetType type, bool async); + + uint getNumTests() { + return _numSubTests; + } + + uint getNumTests2D() { + return _numSubTests2D; + } + uint getNumTests3D() { + return _numSubTests3D; + } +}; + +bool hipPerfMemset::open(int deviceId) { + int nGpu = 0; + HIP_CHECK(hipGetDeviceCount(&nGpu)); + if (nGpu < 1) { + return false; + } + + HIP_CHECK(hipSetDevice(deviceId)); + hipDeviceProp_t props = {0}; + HIP_CHECK(hipGetDeviceProperties(&props, deviceId)); + INFO("info: running on bus " << "0x" << props.pciBusID << " " << props.name + << " with " << props.multiProcessorCount << " CUs and device id: " + << deviceId << "\n"); + return true; +} + +template +void hipPerfMemset::run1D(unsigned int test, T memsetval, + enum MemsetType type, bool async) { + T *A_h, *A_d; + + testTypeSize_ = typeSizeList[(test / num_elements_) % num_typeSize_]; + testNumEle_ = eleNumList[test % num_elements_]; + + bufSize_ = testNumEle_ * 4; + + HIP_CHECK(hipMalloc(&A_d, bufSize_)); + + A_h = reinterpret_cast (malloc(bufSize_)); + + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + // Warm-up + HIP_CHECK(hipMemset(reinterpret_cast(A_d), memsetval, bufSize_)); + + auto start = std::chrono::steady_clock::now(); + + for (uint i = 0; i < NUM_ITER; i++) { + if (type == hipMemsetTypeDefault && !async) { + HIP_CHECK(hipMemset(reinterpret_cast(A_d), memsetval, bufSize_)); + } else if (type == hipMemsetTypeDefault && async) { + HIP_CHECK(hipMemsetAsync(A_d, memsetval, bufSize_, stream)); + } else if (type == hipMemsetTypeD8 && !async) { + HIP_CHECK(hipMemsetD8((hipDeviceptr_t)A_d, memsetval, bufSize_)); + } else if (type == hipMemsetTypeD8 && async) { + HIP_CHECK(hipMemsetD8Async((hipDeviceptr_t)A_d, memsetval, bufSize_)); + } else if (type == hipMemsetTypeD16 && !async) { + HIP_CHECK(hipMemsetD16((hipDeviceptr_t)A_d, memsetval, + bufSize_/sizeof(T))); + } else if (type == hipMemsetTypeD16 && async) { + HIP_CHECK(hipMemsetD16Async((hipDeviceptr_t)A_d, memsetval, + bufSize_/sizeof(T))); + } else if (type == hipMemsetTypeD32 && !async) { + HIP_CHECK(hipMemsetD32((hipDeviceptr_t)A_d, memsetval, + bufSize_/sizeof(T))); + } else if (type == hipMemsetTypeD32 && async) { + HIP_CHECK(hipMemsetD32Async((hipDeviceptr_t)A_d, memsetval, + bufSize_/sizeof(T))); + } + } + + HIP_CHECK(hipDeviceSynchronize()); + + auto end = std::chrono::steady_clock::now(); + + HIP_CHECK(hipMemcpy(A_h, A_d, bufSize_, hipMemcpyDeviceToHost) ); + + for (int i = 0; i < bufSize_/testTypeSize_; i++) { + if (A_h[i] != memsetval) { + INFO("mismatch at index " << i << " computed: " << + static_cast (A_h[i]) << ", memsetval: " << + static_cast (memsetval) << "\n"); + REQUIRE(false); + } + } + + HIP_CHECK(hipFree(A_d)); + free(A_h); + + std::chrono::duration diff = end - start; + + auto sec = diff.count(); + auto perf = static_cast((bufSize_ * NUM_ITER * (1e-09)) / sec); + + INFO("hipPerf1DMemset[" << test << "] " << (int)bufSize_/1024 << " Kb " + << std::setw(4) << " typeSize " << (int) testTypeSize_ << ":" + << std::setw(5) << perf << " GB/s \n"); +} + +template +void hipPerfMemset::run2D(unsigned int test, T memsetval, + enum MemsetType type, bool async) { + bufSize_ = sizeList[test % num_sizes_]; + size_t numH = bufSize_; + size_t numW = bufSize_; + size_t pitch_A; + size_t width = numW * sizeof(char); + size_t sizeElements = width * numH; + size_t elements = numW* numH; + + T * A_h, * A_d; + + HIP_CHECK(hipMallocPitch(reinterpret_cast(&A_d), + &pitch_A, width, numH)); + A_h = reinterpret_cast(malloc(sizeElements)); + + for (size_t i=0; i < elements; i++) { + A_h[i] = 1; + } + + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + // Warm-up + HIP_CHECK(hipMemset2D(A_d, pitch_A, memsetval, numW, numH)); + + auto start = std::chrono::steady_clock::now(); + + for (uint i = 0; i < NUM_ITER; i++) { + if (type == hipMemsetTypeDefault && !async) { + HIP_CHECK(hipMemset2D(A_d, pitch_A, memsetval, numW, numH)); + } else if (type == hipMemsetTypeDefault && async) { + HIP_CHECK(hipMemset2DAsync(A_d, pitch_A, memsetval, numW, numH, stream)); + } + } + + HIP_CHECK(hipStreamSynchronize(stream)); + + auto end = std::chrono::steady_clock::now(); + + HIP_CHECK(hipMemcpy2D(A_h, width, A_d, pitch_A, numW, numH, + hipMemcpyDeviceToHost)); + + for (int i=0; i < elements; i++) { + if (A_h[i] != memsetval) { + INFO("mismatch at index " << i << " computed: " << + static_cast (A_h[i]) << ", memsetval: " << + static_cast (memsetval) << "\n"); + REQUIRE(false); + } + } + + std::chrono::duration diff = end - start; + + auto sec = diff.count(); + auto perf = static_cast((sizeElements* NUM_ITER * (1e-09)) / sec); + + INFO("hipPerf2DMemset[" << test << "] " <<" " << "(GB/s) for " << + (int)bufSize_ << " x " << bufSize_ << " bytes : " << std::setw(5) << + perf << "\n"); + + HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK(hipFree(A_d)); + free(A_h); +} + +template +void hipPerfMemset::run3D(unsigned int test, T memsetval, + enum MemsetType type, bool async) { + bufSize_ = sizeList[test % num_sizes_]; + + size_t numH = bufSize_; + size_t numW = bufSize_; + size_t depth = 10; + size_t width = numW * sizeof(char); + size_t sizeElements = width * numH * depth; + size_t elements = numW* numH* depth; + + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + T *A_h; + + hipExtent extent = make_hipExtent(width, numH, depth); + hipPitchedPtr devPitchedPtr; + + HIP_CHECK(hipMalloc3D(&devPitchedPtr, extent)); + A_h = reinterpret_cast(malloc(sizeElements)); + HIPASSERT(A_h != NULL); + + for (size_t i = 0; i < elements; i++) { + A_h[i] = 1; + } + + // Warm up + HIP_CHECK(hipMemset3D(devPitchedPtr, memsetval, extent)); + + auto start = std::chrono::steady_clock::now(); + + for (uint i = 0; i < NUM_ITER; i++) { + if (type == hipMemsetTypeDefault && !async) { + HIP_CHECK(hipMemset3D(devPitchedPtr, memsetval, extent)); + } else if (type == hipMemsetTypeDefault && async) { + HIP_CHECK(hipMemset3DAsync(devPitchedPtr, memsetval, extent, stream)); + } + } + + HIP_CHECK(hipStreamSynchronize(stream)); + + auto end = std::chrono::steady_clock::now(); + + hipMemcpy3DParms myparms = {0}; + myparms.srcPos = make_hipPos(0, 0, 0); + myparms.dstPos = make_hipPos(0, 0, 0); + myparms.dstPtr = make_hipPitchedPtr(A_h, width , numW, numH); + myparms.srcPtr = devPitchedPtr; + myparms.extent = extent; + + myparms.kind = hipMemcpyDeviceToHost; + + HIP_CHECK(hipMemcpy3D(&myparms)); + + for (int i=0; i < elements; i++) { + if (A_h[i] != memsetval) { + INFO("mismatch at index " << i << " computed: " << + static_cast (A_h[i]) << ", memsetval: " << + static_cast (memsetval) << "\n"); + REQUIRE(false); + } + } + + std::chrono::duration diff = end - start; + + auto sec = diff.count(); + auto perf = static_cast((sizeElements * NUM_ITER * (1e-09)) / sec); + + INFO("hipPerf3DMemset[" << test << "] " <<" " << "(GB/s) for " << + (int)bufSize_ << " x " << bufSize_ << " x " <= 5.6 +*/ + +TEST_CASE("Perf_hipPerfMemset_test") { + hipPerfMemset hipPerfMemset; + + int deviceId = 0; + REQUIRE(hipPerfMemset.open(deviceId)); + + dataType pattern; + + int numTests = hipPerfMemset.getNumTests(); + int numTests2D = hipPerfMemset.getNumTests2D(); + int numTests3D = hipPerfMemset.getNumTests3D(); + + bool async = false; + + for (uint i = 0; i < 2 ; i++) { + if (async) { + INFO("Perf of hipMemsetAsync for 1D arrays \n"); + } else { + INFO("Perf of hipMemset for 1D arrays \n"); + } + + for (auto testCase = 0; testCase < numTests; testCase++) { + if (testCase < 5) { + INFO("API: hipMemset \n"); + hipPerfMemset.run1D(testCase, pattern.memsetval, + hipMemsetTypeDefault, async); + } else if (testCase < 10) { + INFO("API: hipMemsetD16 \n"); + hipPerfMemset.run1D(testCase, pattern.memsetD16val, + hipMemsetTypeD16, async); + } else if (testCase < 15) { + INFO("API: hipMemsetD32 \n"); + hipPerfMemset.run1D(testCase, pattern.memsetD32val, + hipMemsetTypeD32, async); + } else { + INFO("API: hipMemset \n"); + hipPerfMemset.run1D(testCase, pattern.memsetval, + hipMemsetTypeDefault, async); + } + } + async = true; + } + + for (uint i = 0; i < 2; i++) { + if (async) { + INFO("Perf of hipMemset2DAsync for 2D arrays \n"); + } else { + INFO("Perf of hipMemset2D for 2D arrays \n"); + } + + for (uint test = 0; test < numTests2D; test++) { + hipPerfMemset.run2D(test, pattern.memsetval, hipMemsetTypeDefault, async); + } + async = false; + } + + for (uint i = 0; i < 2; i++) { + if (async) { + INFO("Perf of hipMemset3DAsync for 3D arrays \n"); + } else { + INFO("Perf of hipMemset3D for 3D arrays \n"); + } + + for (uint test = 0; test < numTests3D; test++) { + hipPerfMemset.run3D(test, pattern.memsetval, hipMemsetTypeDefault, async); + } + } +} diff --git a/projects/hip-tests/catch/perftests/memory/hipPerfSampleRate.cc b/projects/hip-tests/catch/perftests/memory/hipPerfSampleRate.cc new file mode 100644 index 0000000000..1f5ac59ffa --- /dev/null +++ b/projects/hip-tests/catch/perftests/memory/hipPerfSampleRate.cc @@ -0,0 +1,311 @@ +/* + Copyright (c) 2023 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. + */ + + +/** +* @addtogroup hipMemcpyKernel hipMemcpyKernel +* @{ +* @ingroup perfMemoryTest +* `hipMemcpy(void* dst, const void* src, size_t count, hipMemcpyKind kind)` - +* Copies data between host and device. +*/ + +#include + +#define NUM_TYPES 3 +std::vector types = {"float", "float2", "float4"}; +std::vector typeSizes = {4, 8, 16}; + +#define NUM_SIZES 12 +std::vector sizes = {1, 2, 4, 8, 16, 32, + 64, 128, 256, 512, 1024, 2048}; + +#define NUM_BUFS 6 +#define MAX_BUFS (1 << (NUM_BUFS - 1)) + +#ifdef __HIP_PLATFORM_NVIDIA__ +__host__ __device__ void operator+=(float2 &a, float2 b) { //NOLINT + a.x += b.x; a.y += b.y; +} + +__host__ __device__ void operator+=(float4 &a, float4 b) { //NOLINT + a.x += b.x; a.y += b.y; a.z += b.z; a.w += b.w; +} +#endif + +template +__global__ void sampleRate(T * outBuffer, unsigned int inBufSize, + unsigned int writeIt, T **inBuffer, int numBufs) { + uint gid = (blockIdx.x * blockDim.x + threadIdx.x); + uint inputIdx = gid % inBufSize; + + T tmp; + memset(&tmp, 0, sizeof(T)); + for (int i = 0; i < numBufs; i++) { + tmp += *(*(inBuffer+i)+inputIdx); + } + + if (writeIt*(unsigned int)tmp.x) { + outBuffer[gid] = tmp; + } +} + +template +__global__ void sampleRateFloat(T * outBuffer, unsigned int inBufSize, + unsigned int writeIt, T ** inBuffer, int numBufs) { + uint gid = (blockIdx.x * blockDim.x + threadIdx.x); + uint inputIdx = gid % inBufSize; + + T tmp = (T)0.0f; + + for (int i = 0; i < numBufs; i++) { + tmp += *((*inBuffer+i)+inputIdx); + } + + if (writeIt*(unsigned int)tmp) { + outBuffer[gid] = tmp; + } +} + +class hipPerfSampleRate { + public: + hipPerfSampleRate(); + ~hipPerfSampleRate(); + + bool open(void); + void run(unsigned int testCase); + void close(void); + + // array of funtion pointers + typedef void (hipPerfSampleRate::*funPtr)(void * outBuffer, unsigned int + inBufSize, unsigned int writeIt, void **inBuffer, int numBufs, + int grids, int blocks, int threads_per_block); + + // Wrappers + void float_kernel(void * outBuffer, unsigned int inBufSize, + unsigned int writeIt, void **inBuffer, int numBufs, + int grids, int blocks, int threads_per_block); + + void float2_kernel(void * outBuffer, unsigned int inBufSize, + unsigned int writeIt, void **inBuffer, int numBufs, + int grids, int blocks, int threads_per_block); + + void float4_kernel(void * outBuffer, unsigned int inBufSize, + unsigned int writeIt, void **inBuffer, int numBufs, + int grids, int blocks, int threads_per_block); + + private: + void setData(void *ptr, unsigned int value); + void checkData(uint *ptr); + + unsigned int width_; + unsigned int bufSize_; + int numCUs; + + unsigned int outBufSize_; + static const unsigned int MAX_ITERATIONS = 25; + unsigned int numBufs_; + unsigned int typeIdx_; +}; + +hipPerfSampleRate::hipPerfSampleRate() {} +hipPerfSampleRate::~hipPerfSampleRate() {} +void hipPerfSampleRate::close() {} + +bool hipPerfSampleRate::open(void) { + int nGpu = 0; + HIP_CHECK(hipGetDeviceCount(&nGpu)); + if (nGpu < 1) { + return false; + } + + int deviceId = 0; + hipDeviceProp_t props = {0}; + props = {0}; + HIP_CHECK(hipSetDevice(deviceId)); + HIP_CHECK(hipGetDeviceProperties(&props, deviceId)); + INFO("info: running on bus " << "0x" << props.pciBusID << " " << + props.name << " with " << props.multiProcessorCount << + " CUs" << " and device id: " << deviceId << "\n"); + numCUs = props.multiProcessorCount; + return true; +} + +// Wrappers for the kernel launches +void hipPerfSampleRate::float_kernel(void * outBuffer, unsigned int inBufSize, + unsigned int writeIt, void **inBuffer, int numBufs, + int grids, int blocks, int threads_per_block) { + hipLaunchKernelGGL(sampleRateFloat, dim3(grids, grids, grids), + dim3(blocks), 0, 0, reinterpret_cast(outBuffer), + inBufSize, writeIt, reinterpret_cast(inBuffer), numBufs); +} + +void hipPerfSampleRate::float2_kernel(void * outBuffer, unsigned int inBufSize, + unsigned int writeIt, void **inBuffer, int grids, + int blocks, int threads_per_block, int numBufs) { + hipLaunchKernelGGL(sampleRate, dim3(grids, grids, grids), + dim3(blocks), 0, 0, reinterpret_cast(outBuffer), + inBufSize, writeIt, reinterpret_cast(inBuffer), numBufs); +} + +void hipPerfSampleRate::float4_kernel(void * outBuffer, unsigned int inBufSize, + unsigned int writeIt, void **inBuffer, int grids, + int blocks, int threads_per_block, int numBufs) { + hipLaunchKernelGGL(sampleRate, dim3(grids, grids, grids), + dim3(blocks), 0, 0, reinterpret_cast(outBuffer), + inBufSize, writeIt, reinterpret_cast(inBuffer), numBufs); +} + +void hipPerfSampleRate::run(unsigned int test) { + funPtr p[] = {&hipPerfSampleRate::float_kernel, + &hipPerfSampleRate::float2_kernel, + &hipPerfSampleRate::float4_kernel}; + + // We compute a square domain + width_ = sizes[test % NUM_SIZES]; + typeIdx_ = (test / NUM_SIZES) % NUM_TYPES; + bufSize_ = width_ * width_ * typeSizes[typeIdx_]; + numBufs_ = (1 << (test / (NUM_SIZES * NUM_TYPES))); + + void ** dPtr; + void * hOutPtr; + void * dOutPtr; + void * hInPtr[numBufs_]; + void * dInPtr[numBufs_]; + + outBufSize_ = + sizes[NUM_SIZES - 1] * sizes[NUM_SIZES - 1] * typeSizes[NUM_TYPES - 1]; + + // Allocate memory on the host and device + HIP_CHECK(hipHostMalloc(reinterpret_cast(&hOutPtr), outBufSize_, + hipHostMallocDefault)); + setData(reinterpret_cast(hOutPtr), 0xdeadbeef); + HIP_CHECK(hipMalloc(reinterpret_cast(&dOutPtr), outBufSize_)); + + // Allocate 2D array in Device + HIP_CHECK(hipMalloc(reinterpret_cast(&dPtr), + numBufs_* sizeof(void *))); + + for (uint i = 0; i < numBufs_; i++) { + HIP_CHECK(hipHostMalloc(reinterpret_cast(&hInPtr[i]), bufSize_, + hipHostMallocDefault)); + HIP_CHECK(hipMalloc(reinterpret_cast(&dInPtr[i]), bufSize_)); + setData(hInPtr[i], 0x3f800000); + } + + // Populate array of pointers with array addresses + HIP_CHECK(hipMemcpy(dPtr, dInPtr, numBufs_* sizeof(void *), + hipMemcpyHostToDevice)); + + // Copy memory from host to device + for (uint i = 0; i < numBufs_; i++) { + HIP_CHECK(hipMemcpy(dInPtr[i], hInPtr[i], bufSize_, hipMemcpyHostToDevice)); + } + + HIP_CHECK(hipMemcpy(dOutPtr, hOutPtr, outBufSize_, hipMemcpyHostToDevice)); + + // Prepare kernel launch parameters + // outBufSize_/sizeof(uint) - Grid size in 3D + int grids = 64; + int blocks = 64; + int threads_per_block = 1; + + unsigned int maxIter = MAX_ITERATIONS * (MAX_BUFS / numBufs_); + unsigned int sizeDW = width_ * width_; + unsigned int writeIt = 0; + + int idx = 0; + + if (!types[typeIdx_].compare("float")) { + idx = 0; + } else if (!types[typeIdx_].compare("float2")) { + idx = 1; + } else if (!types[typeIdx_].compare("float4")) { + idx = 2; + } + + // Time the kernel execution + auto all_start = std::chrono::steady_clock::now(); + for (uint i = 0; i < maxIter; i++) { + (this->*p[idx]) (reinterpret_cast(dOutPtr), sizeDW, writeIt, + dPtr, numBufs_, grids, blocks, threads_per_block); + } + + HIP_CHECK(hipDeviceSynchronize()); + auto all_end = std::chrono::steady_clock::now(); + std::chrono::duration all_kernel_time = all_end - all_start; + + double perf = (static_cast(outBufSize_ * numBufs_ * + maxIter * (1e-09))) / all_kernel_time.count(); + + INFO("Domain " << sizes[NUM_SIZES - 1] << "x"<< sizes[NUM_SIZES - 1] + << " bufs " << numBufs_ << " " << types[typeIdx_] << " " << width_ + << "x" <(numBufs_)) { + INFO("Data validation failed at "<< i << " Got "<< ptr[i] + << ", expected " << (float)numBufs_ << "\n"); + REQUIRE(false); + } + } +} + +/** +* Test Description +* ------------------------ +*  - Verify hipPerfSampleRate status. +* Test source +* ------------------------ +*  - perftests/memory/hipPerfSampleRate.cc +* Test requirements +* ------------------------ +*  - HIP_VERSION >= 5.6 +*/ + +TEST_CASE("Perf_hipPerfSampleRate_test") { + hipPerfSampleRate sampleTypes; + + REQUIRE(true == sampleTypes.open()); + + for (unsigned int testCase = 0; testCase < 216 ; testCase+=36) { + sampleTypes.run(testCase); + } +} diff --git a/projects/hip-tests/catch/perftests/memory/hipPerfSharedMemReadSpeed.cc b/projects/hip-tests/catch/perftests/memory/hipPerfSharedMemReadSpeed.cc new file mode 100644 index 0000000000..1d4b620c0e --- /dev/null +++ b/projects/hip-tests/catch/perftests/memory/hipPerfSharedMemReadSpeed.cc @@ -0,0 +1,262 @@ +/* + Copyright (c) 2023 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. + */ + +/** +* @addtogroup hipMemcpyKernel hipMemcpyKernel +* @{ +* @ingroup perfMemoryTest +* `hipMemcpy(void* dst, const void* src, size_t count, hipMemcpyKind kind)` - +* Copies data between host and device. +*/ + +#include + +#define sharedMemSize1 2048 +#define sharedMemSize2 256 + +__global__ void sharedMemReadSpeed1(float *outBuf, ulong N) { + size_t gid = (blockIdx.x * blockDim.x + threadIdx.x); + size_t lid = threadIdx.x; + __shared__ float local[sharedMemSize1]; + + float val1 = 0; + float val2 = 0; + float val3 = 0; + float val4 = 0; + + for (int i = 0; i < (sharedMemSize1 / 64); i++) { + local[lid + i * 64] = lid; + } + + __syncthreads(); + + val1 += local[lid]; + val2 += local[lid + 64]; + val3 += local[lid + 128]; + val4 += local[lid + 192]; + val1 += local[lid + 256]; + val2 += local[lid + 320]; + val3 += local[lid + 384]; + val4 += local[lid + 448]; + val1 += local[lid + 512]; + val2 += local[lid + 576]; + val3 += local[lid + 640]; + val4 += local[lid + 704]; + val1 += local[lid + 768]; + val2 += local[lid + 832]; + val3 += local[lid + 896]; + val4 += local[lid + 960]; + val1 += local[lid + 1024]; + val2 += local[lid + 1088]; + val3 += local[lid + 1152]; + val4 += local[lid + 1216]; + val1 += local[lid + 1280]; + val2 += local[lid + 1344]; + val3 += local[lid + 1408]; + val4 += local[lid + 1472]; + val1 += local[lid + 1536]; + val2 += local[lid + 1600]; + val3 += local[lid + 1664]; + val4 += local[lid + 1728]; + val1 += local[lid + 1792]; + val2 += local[lid + 1856]; + val3 += local[lid + 1920]; + val4 += local[lid + 1984]; + + if (gid < N) { + outBuf[gid] = val1 + val2 + val3 + val4; + } +} + +__global__ void sharedMemReadSpeed2(float *outBuf, ulong N) { + size_t gid = (blockIdx.x * blockDim.x + threadIdx.x); + size_t lid = threadIdx.x; + __shared__ float local[sharedMemSize2]; + + float val0 = 0.0f; + float val1 = 0.0f; + + for (int i = 0; i < (sharedMemSize2 / 64); i++) { + local[lid + i * 64] = lid; + } + + __syncthreads(); + +#pragma nounroll + for (uint i = 0; i < 32; i++) { + val0 += local[8 * i + 0]; + val1 += local[8 * i + 1]; + val0 += local[8 * i + 2]; + val1 += local[8 * i + 3]; + val0 += local[8 * i + 4]; + val1 += local[8 * i + 5]; + val0 += local[8 * i + 6]; + val1 += local[8 * i + 7]; + } + + if (gid < N) { + outBuf[gid] = val0 + val1; + } +} + +static bool hipPerfSharedMemReadSpeed_test() { + float *dDst; + float *hDst; + hipStream_t stream; + constexpr uint numSizes = 4; + constexpr uint Sizes[numSizes] = {262144, 1048576, 4194304, 16777216}; + uint numReads1 = 32; + uint numReads2 = 256; + uint sharedMemSizeBytes1 = sharedMemSize1 * sizeof(float); + uint sharedMemSizeBytes2 = sharedMemSize2 * sizeof(float); + int nIter = 1000; + const unsigned threadsPerBlock = 64; + + static int device = 0; + HIP_CHECK(hipSetDevice(device)); + hipDeviceProp_t props; + HIP_CHECK(hipGetDeviceProperties(&props, device)); + INFO("info: running on bus " << "0x" << props.pciBusID << " " << props.name + << " with " << props.multiProcessorCount << " CUs \n"); + + HIP_CHECK(hipStreamCreate(&stream)); + + for (int nTest = 0; nTest < numSizes; nTest++) { + uint nBytes = Sizes[nTest % numSizes]; + ulong N = nBytes / sizeof(float); + const unsigned blocks = N / threadsPerBlock; + + hDst = new float[nBytes]; + HIP_CHECK(hDst == 0 ? hipErrorOutOfMemory : hipSuccess); + memset(hDst, 0, nBytes); + + HIP_CHECK(hipMalloc(&dDst, nBytes)); + HIP_CHECK(hipMemcpy(dDst, hDst, nBytes, hipMemcpyHostToDevice)); + + hipLaunchKernelGGL(sharedMemReadSpeed1, dim3(blocks), + dim3(threadsPerBlock), 0, stream, dDst, N); + HIP_CHECK(hipMemcpy(hDst, dDst, nBytes, hipMemcpyDeviceToHost)); + HIP_CHECK(hipDeviceSynchronize()); + + int tmp = 0; + for (int i = 0; i < N; i++) { + if (i % threadsPerBlock == 0) { + tmp = 0; + } + if (hDst[i] != tmp) { + INFO("info: Data validation failed for warm up run! \n"); + INFO("info: expected " << tmp << " got " << hDst[i] << " \n"); + return false; + } + tmp += threadsPerBlock / 2; + } + + auto all_start = std::chrono::steady_clock::now(); + for (int i = 0; i < nIter; i++) { + hipLaunchKernelGGL(sharedMemReadSpeed1, dim3(blocks), + dim3(threadsPerBlock), 0, stream, dDst, N); + } + HIP_CHECK(hipDeviceSynchronize()); + + auto all_end = std::chrono::steady_clock::now(); + std::chrono::duration all_kernel_time = all_end - all_start; + + // read speed in GB/s + double perf = (static_cast(blocks * threadsPerBlock) + * (numReads1 * sizeof(float) + sharedMemSizeBytes1 / 64) + * nIter * (1e-09)) / all_kernel_time.count(); + + INFO("info: read speed = " << std::setw(8) << perf << " GB/s for " << + sharedMemSizeBytes1 / 1024 << " KB shared memory with " << + std::setw(8) << blocks * threadsPerBlock << " threads, " + << std::setw(4) << numReads1 << + " reads in sharedMemReadSpeed1 kernel \n"); + + delete[] hDst; + HIP_CHECK(hipFree(dDst)); + } + + for (int nTest = 0; nTest < numSizes; nTest++) { + uint nBytes = Sizes[nTest % numSizes]; + ulong N = nBytes / sizeof(float); + const unsigned blocks = N / threadsPerBlock; + + hDst = new float[nBytes]; + HIP_CHECK(hDst == 0 ? hipErrorOutOfMemory : hipSuccess); + memset(hDst, 0, nBytes); + + HIP_CHECK(hipMalloc(&dDst, nBytes)); + HIP_CHECK(hipMemcpy(dDst, hDst, nBytes, hipMemcpyHostToDevice)); + + hipLaunchKernelGGL(sharedMemReadSpeed2, dim3(blocks), + dim3(threadsPerBlock), 0, stream, dDst, N); + HIP_CHECK(hipMemcpy(hDst, dDst, nBytes, hipMemcpyDeviceToHost)); + HIP_CHECK(hipDeviceSynchronize()); + + auto all_start = std::chrono::steady_clock::now(); + for (int i = 0; i < nIter; i++) { + hipLaunchKernelGGL(sharedMemReadSpeed2, dim3(blocks), + dim3(threadsPerBlock), 0, stream, dDst, N); + } + HIP_CHECK(hipDeviceSynchronize()); + + auto all_end = std::chrono::steady_clock::now(); + std::chrono::duration all_kernel_time = all_end - all_start; + + // read speed in GB/s + double perf = (static_cast(blocks * threadsPerBlock) + * (numReads2 * sizeof(float) + sharedMemSizeBytes2 / 64) + * nIter * (1e-09)) / all_kernel_time.count(); + + INFO("info: read speed = " << std::setw(8) << perf << " GB/s for " + << sharedMemSizeBytes2 / 1024 << " KB shared memory with " + << std::setw(8) << blocks * threadsPerBlock << " threads, " + << std::setw(4) << numReads2 << + " reads in sharedMemReadSpeed2 kernel \n"); + + delete[] hDst; + HIP_CHECK(hipFree(dDst)); + } + HIP_CHECK(hipStreamDestroy(stream)); + return true; +} + +/** +* Test Description +* ------------------------ +*  - Verify hipPerfSharedMemReadSpeed status. +* Test source +* ------------------------ +*  - perftests/memory/hipPerfSharedMemReadSpeed.cc +* Test requirements +* ------------------------ +*  - HIP_VERSION >= 5.6 +*/ + +TEST_CASE("Perf_hipPerfSharedMemReadSpeed_test") { + int numDevices = 0; + HIP_CHECK(hipGetDeviceCount(&numDevices)); + + if (numDevices <= 0) { + SUCCEED("Skipped testcase hipPerfSharedMemReadSpeed as" + "there is no device to test.\n"); + } else { + REQUIRE(true == hipPerfSharedMemReadSpeed_test()); + } +}