diff --git a/projects/hip-tests/catch/perftests/CMakeLists.txt b/projects/hip-tests/catch/perftests/CMakeLists.txt index 35751d79c1..38078c0a48 100644 --- a/projects/hip-tests/catch/perftests/CMakeLists.txt +++ b/projects/hip-tests/catch/perftests/CMakeLists.txt @@ -26,3 +26,4 @@ add_subdirectory(stream) add_subdirectory(dispatch) add_subdirectory(compute) add_subdirectory(graph) +add_subdirectory(event) diff --git a/projects/hip-tests/catch/perftests/event/CMakeLists.txt b/projects/hip-tests/catch/perftests/event/CMakeLists.txt new file mode 100644 index 0000000000..8d233224fa --- /dev/null +++ b/projects/hip-tests/catch/perftests/event/CMakeLists.txt @@ -0,0 +1,24 @@ +# Copyright (c) 2025 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. +# Common Tests - Test independent of all platforms +set(TEST_SRC + hipKernelLookUpPerf.cc +) +hip_add_exe_to_target(NAME perfEventTest + TEST_SRC ${TEST_SRC} + TEST_TARGET_NAME perf_test) + diff --git a/projects/hip-tests/catch/perftests/event/hipKernelLookUpPerf.cc b/projects/hip-tests/catch/perftests/event/hipKernelLookUpPerf.cc new file mode 100644 index 0000000000..f237c219f2 --- /dev/null +++ b/projects/hip-tests/catch/perftests/event/hipKernelLookUpPerf.cc @@ -0,0 +1,134 @@ +/*Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ +#include +#include +#include +#include +#include +#include +#define HIP_CHECK_PERF(a) \ + { \ + auto err = a; \ + if ((err != hipSuccess) && (err != hipErrorNotReady)) { \ + printf(#a "= Error! %s\n", hipGetErrorString(err)); \ + exit(1); \ + } \ + } +/** + * @addtogroup hipLaunchKernelGGL hipLaunchKernelGGL + * @{ + * @ingroup PerformanceTest + */ +__global__ void empty_kernel() { + __shared__ int temp[256]; + temp[threadIdx.x] = sinf(float(threadIdx.x)); +} +void rocm_empty_gpu_job(void *stream) { + hipLaunchKernelGGL(empty_kernel, 1, 256, 0, (hipStream_t)stream); +} +std::vector> stream_pools; +std::atomic count(0); +bool kill = false; +std::chrono::system_clock::time_point thread_report[16]; +void thread_jobs(int dev, int virt) { + HIP_CHECK_PERF(hipSetDevice(dev)); + hipStream_t exec_stream = stream_pools[dev][virt]; + uint64_t n = 0; + while (!kill) { + rocm_empty_gpu_job(exec_stream); + n++; + if ((n & 150) == 0) { + HIP_CHECK_PERF(hipStreamSynchronize(exec_stream)); + thread_report[dev * 4 + virt] = std::chrono::system_clock::now(); + } + count++; + } + HIP_CHECK_PERF(hipStreamSynchronize(exec_stream)); +} +/** + * Test Description + * ------------------------ + * - This test case prints the number of jobs/Second. + * - 1) Launch number of threads on each device. + * - 2) Threads should have only dispatches. + * - 3) In the main thread calculate the number of jobs/Second + * - 4) Print the jobs/Second value. + * Test source + * ------------------------ + * - catch/perftests/event/hipKernelLookUpPerf.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.5 + */ +TEST_CASE("Unit_hipKernelLookUp_PerfTest") { + int mgpu = 0; + HIP_CHECK_PERF(hipGetDeviceCount(&mgpu)); + stream_pools.resize(mgpu); + HIP_CHECK_PERF(hipSetDeviceFlags(hipDeviceScheduleSpin)); + for (int i = 0; i < mgpu; i++) { + HIP_CHECK_PERF(hipSetDevice(i)); + stream_pools[i].resize(12); + for (int j = 0; j < 12; j++) + HIP_CHECK_PERF( + hipStreamCreateWithFlags(&stream_pools[i][j], hipStreamNonBlocking)); + } + for (int nDev = 1; nDev <= mgpu; nDev++) { + count = 0; + INFO("RUNNING ON "< threads; + for (int i = 0; i < nDev * 4; i++) + threads.push_back(std::thread(thread_jobs, i / 4, i % 4)); + usleep(1000000); + auto t1 = std::chrono::system_clock::now(); + int counter = int(count); + uint64_t total_count = 0; + double total_time = 0; + for (int t = 0; t < 10; t++) { + usleep(1000000); + auto t2 = std::chrono::system_clock::now(); + auto duration = + std::chrono::duration_cast(t2 - t1) + .count(); + int counter2 = int(count); + for (int i = 0; i < nDev * 4; i++) { + if (std::chrono::duration_cast( + t2 - thread_report[i]) + .count() >= 1000000) { + INFO("Thread "<