From 72e64327397e5c1d6f13a4439b303052d5ff1ff2 Mon Sep 17 00:00:00 2001 From: milos-mozetic <118800401+milos-mozetic@users.noreply.github.com> Date: Mon, 6 Mar 2023 04:28:26 +0100 Subject: [PATCH] EXSWHTEC-174 - Implement unit tests for Callback Activity APIs (#6) Implemented positive and negative test cases for the following APIs: - hipApiName - hipKernelNameRef - hipKernelNameRefByPtr - hipGetStreamDeviceId Enabled build of Callback tests only for AMD platforms, as there is no support for Nvidia platforms. --- .../config/config_amd_linux_common.json | 1 + .../config/config_amd_windows_common.json | 2 + catch/unit/CMakeLists.txt | 1 + catch/unit/callback/CMakeLists.txt | 44 +++++++ catch/unit/callback/SimpleKernel.cc | 27 +++++ catch/unit/callback/hipApiName.cc | 79 ++++++++++++ catch/unit/callback/hipGetStreamDeviceId.cc | 112 ++++++++++++++++++ catch/unit/callback/hipKernelNameRef.cc | 71 +++++++++++ catch/unit/callback/hipKernelNameRefByPtr.cc | 105 ++++++++++++++++ 9 files changed, 442 insertions(+) create mode 100644 catch/unit/callback/CMakeLists.txt create mode 100644 catch/unit/callback/SimpleKernel.cc create mode 100644 catch/unit/callback/hipApiName.cc create mode 100644 catch/unit/callback/hipGetStreamDeviceId.cc create mode 100644 catch/unit/callback/hipKernelNameRef.cc create mode 100644 catch/unit/callback/hipKernelNameRefByPtr.cc diff --git a/catch/hipTestMain/config/config_amd_linux_common.json b/catch/hipTestMain/config/config_amd_linux_common.json index a4e1fd0cb1..5fdaefd781 100644 --- a/catch/hipTestMain/config/config_amd_linux_common.json +++ b/catch/hipTestMain/config/config_amd_linux_common.json @@ -15,6 +15,7 @@ "Unit_hipMemset_Negative_OutOfBoundsPtr", "Unit_hipDeviceReset_Positive_Basic", "Unit_hipDeviceReset_Positive_Threaded", + "Unit_hipKernelNameRef_Negative_Parameters", "Unit_hipMemAdvise_AccessedBy_All_Devices", "Unit_hipMemAdvise_No_Flag_Interference", "Unit_hipGraphDestroyNode_Complx_ChkNumOfNodesNDep", diff --git a/catch/hipTestMain/config/config_amd_windows_common.json b/catch/hipTestMain/config/config_amd_windows_common.json index 73265230ea..5bfc0356d7 100644 --- a/catch/hipTestMain/config/config_amd_windows_common.json +++ b/catch/hipTestMain/config/config_amd_windows_common.json @@ -95,6 +95,8 @@ "Unit_hipStreamSynchronize_NullStreamAndStreamPerThread", "Note: intermittent Seg fault failure ", "Unit_hipGraphAddEventRecordNode_Functional_WithoutFlags", + "Unit_hipKernelNameRef_Negative_Parameters", + "Unit_hipKernelNameRef_Positive_Basic", "Unit_hipMemAdvise_AccessedBy_All_Devices", "Unit_hipMemAdvise_No_Flag_Interference", "Unit_hipGraphAddEventRecordNode_Functional_WithoutFlags", diff --git a/catch/unit/CMakeLists.txt b/catch/unit/CMakeLists.txt index 57eb632947..63bbeda241 100644 --- a/catch/unit/CMakeLists.txt +++ b/catch/unit/CMakeLists.txt @@ -35,6 +35,7 @@ add_subdirectory(compiler) add_subdirectory(errorHandling) add_subdirectory(cooperativeGrps) if(HIP_PLATFORM STREQUAL "amd") +add_subdirectory(callback) #add_subdirectory(clock) # Vulkan interop APIs currently undefined for Nvidia add_subdirectory(vulkan_interop) diff --git a/catch/unit/callback/CMakeLists.txt b/catch/unit/callback/CMakeLists.txt new file mode 100644 index 0000000000..204d5bc32c --- /dev/null +++ b/catch/unit/callback/CMakeLists.txt @@ -0,0 +1,44 @@ +# Copyright (c) 2022 Advanced Micro Devices, Inc. All Rights Reserved. +# +# Permission is hereby granted, free of charge, to any person obtaining a copy +# of this software and associated documentation files (the "Software"), to deal +# in the Software without restriction, including without limitation the rights +# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +# copies of the Software, and to permit persons to whom the Software is +# furnished to do so, subject to the following conditions: +# +# The above copyright notice and this permission notice shall be included in +# all copies or substantial portions of the Software. +# +# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +# THE SOFTWARE. + +# Common Tests - Test independent of all platforms + +set(TEST_SRC + hipApiName.cc + hipGetStreamDeviceId.cc + hipKernelNameRef.cc +) + +if(UNIX) + set(TEST_SRC ${TEST_SRC} hipKernelNameRefByPtr.cc) + + add_custom_target(SimpleKernel.code COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR} + ${CMAKE_CURRENT_SOURCE_DIR}/SimpleKernel.cc + -o ${CMAKE_CURRENT_BINARY_DIR}/../../unit/callback/SimpleKernel.code + -I${HIP_PATH}/include -I${CMAKE_CURRENT_SOURCE_DIR}/../../include) +endif() + +hip_add_exe_to_target(NAME CallbackTest + TEST_SRC ${TEST_SRC} + TEST_TARGET_NAME build_tests) + +if(UNIX) + add_dependencies(CallbackTest SimpleKernel.code) +endif() diff --git a/catch/unit/callback/SimpleKernel.cc b/catch/unit/callback/SimpleKernel.cc new file mode 100644 index 0000000000..3bb045ffd1 --- /dev/null +++ b/catch/unit/callback/SimpleKernel.cc @@ -0,0 +1,27 @@ +/* +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include "hip/hip_runtime.h" + +extern "C" __global__ void simple_kernel() { + printf("Hello World!"); +} diff --git a/catch/unit/callback/hipApiName.cc b/catch/unit/callback/hipApiName.cc new file mode 100644 index 0000000000..51f778c4dc --- /dev/null +++ b/catch/unit/callback/hipApiName.cc @@ -0,0 +1,79 @@ +/* +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include + +/** + * @addtogroup hipApiName hipApiName + * @{ + * @ingroup CallbackTest + * `hipApiName(uint32_t id)` - + * returns the name of API with passed ID + */ + +const char* kUnknownApi{"unknown"}; +const uint32_t kApiNumber{1024}; + +/** + * Test Description + * ------------------------ + * - Acquires HIP API names and checks that they are valid + * Test source + * ------------------------ + * - unit/callback/hipApiName.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + * - Platform specific (AMD) + */ +TEST_CASE("Unit_hipApiName_Positive_Basic") { + std::vector hip_api_names; + + for(uint32_t i = 0; i < kApiNumber; ++i) { + if(strcmp(hipApiName(i), kUnknownApi)) { + hip_api_names.emplace_back(hipApiName(i)); + } + } + + REQUIRE(!hip_api_names.empty()); +} + +/** + * Test Description + * ------------------------ + * - Checks that upper and lower limit IDs are mapped to unknown APIs: + * -# When the `uint32_t` upper limit is passed + * - Expected output: return "unknown" + * -# When the `uint32_t` lower limit is passed + * - Expected output: return "unknown" + * Test source + * ------------------------ + * - unit/callback/hipApiName.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + * - Platform specific (AMD) + */ +TEST_CASE("Unit_hipApiName_Negative_ReservedIds") { + REQUIRE_THAT(hipApiName(std::numeric_limits::min()), Catch::Equals(kUnknownApi)); + REQUIRE_THAT(hipApiName(std::numeric_limits::max()), Catch::Equals(kUnknownApi)); +} diff --git a/catch/unit/callback/hipGetStreamDeviceId.cc b/catch/unit/callback/hipGetStreamDeviceId.cc new file mode 100644 index 0000000000..e8d718387e --- /dev/null +++ b/catch/unit/callback/hipGetStreamDeviceId.cc @@ -0,0 +1,112 @@ +/* +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include + +/** + * @addtogroup hipGetStreamDeviceId hipGetStreamDeviceId + * @{ + * @ingroup CallbackTest + * `hipGetStreamDeviceId(hipStream_t stream)` - + * returns the ID of the device on which the stream is active + */ + +/** + * Test Description + * ------------------------ + * - Creates a new stream for each available device + * - Verifies that the Device Stream ID is equal to the Device ID + * Test source + * ------------------------ + * - unit/callback/hipGetStreamDeviceId.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + * - Platform specific (AMD) + */ +TEST_CASE("Unit_hipGetStreamDeviceId_Positive_Threaded_Basic") { + int id = GENERATE(range(0, HipTest::getDeviceCount())); + HIP_CHECK(hipSetDevice(id)); + + StreamGuard stream_guard{Streams::created}; + REQUIRE(hipGetStreamDeviceId(stream_guard.stream()) == id); +} + +/** + * Test Description + * ------------------------ + * - Creates a new stream for each available device, through multiple threads + * - Verifies that the Device Stream ID is equal to the Device ID, from each thread + * Test source + * ------------------------ + * - unit/callback/hipGetStreamDeviceId.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + * - Platform specific (AMD) + * - Multithreaded GPU + */ +TEST_CASE("Unit_hipGetStreamDeviceId_Positive_Multithreaded_Basic") { + const unsigned int max_threads = std::thread::hardware_concurrency(); + const int device_count = HipTest::getDeviceCount(); + + auto thread_function = [&]() { + for(unsigned int id = 0; id < device_count; ++id) { + HIP_CHECK_THREAD(hipSetDevice(id)); + + StreamGuard stream_guard{Streams::perThread}; + REQUIRE_THREAD(hipGetStreamDeviceId(stream_guard.stream()) == id); + } + }; + + std::vector thread_pool; + for(unsigned int i = 0; i < max_threads; ++i) { + thread_pool.emplace_back(thread_function); + } + + for(auto& thread: thread_pool) + { + thread.join(); + } + + HIP_CHECK_THREAD_FINALIZE(); +} + +/** + * Test Description + * ------------------------ + * - Checks that function returns valid ID if the stream is `nullptr` + * Test source + * ------------------------ + * - unit/callback/hipGetStreamDeviceId.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + * - Platform specific (AMD) + */ +TEST_CASE("Unit_hipGetStreamDeviceId_Negative_Parameters") { + int id = GENERATE(range(0, HipTest::getDeviceCount())); + HIP_CHECK(hipSetDevice(id)); + + StreamGuard stream_guard{Streams::nullstream}; + REQUIRE(hipGetStreamDeviceId(stream_guard.stream()) == id); +} diff --git a/catch/unit/callback/hipKernelNameRef.cc b/catch/unit/callback/hipKernelNameRef.cc new file mode 100644 index 0000000000..ff9364c3db --- /dev/null +++ b/catch/unit/callback/hipKernelNameRef.cc @@ -0,0 +1,71 @@ +/* +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include + +/** + * @addtogroup hipKernelNameRef hipKernelNameRef + * @{ + * @ingroup CallbackTest + * `hipKernelNameRef(const hipFunction_t f)` - + * returns the name of passed function object + */ + +/** + * Test Description + * ------------------------ + * - Loads the simple kernel function from the matching module + * - Checks that the valid name is returned for the loaded kernel function + * Test source + * ------------------------ + * - unit/callback/hipKernelNameRef.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + * - Platform specific (AMD) + */ +TEST_CASE("Unit_hipKernelNameRef_Positive_Basic") { + hipModule_t kernel_module{nullptr}; + hipFunction_t kernel_function{nullptr}; + + HIP_CHECK(hipModuleLoad(&kernel_module, "SimpleKernel.code")); + HIP_CHECK(hipModuleGetFunction(&kernel_function, kernel_module, "simple_kernel")); + REQUIRE(hipKernelNameRef(kernel_function) != nullptr); + HIP_CHECK(hipModuleUnload(kernel_module)); +} + +/** + * Test Description + * ------------------------ + * - Checks that the API returns nullptr if the passed function is not loaded + * Test source + * ------------------------ + * - unit/callback/hipKernelNameRef.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + * - Platform specific (AMD) + */ +TEST_CASE("Unit_hipKernelNameRef_Negative_Parameters") { + hipFunction_t kernel_function{nullptr}; + REQUIRE(hipKernelNameRef(kernel_function) == nullptr); +} diff --git a/catch/unit/callback/hipKernelNameRefByPtr.cc b/catch/unit/callback/hipKernelNameRefByPtr.cc new file mode 100644 index 0000000000..dc4d7205f4 --- /dev/null +++ b/catch/unit/callback/hipKernelNameRefByPtr.cc @@ -0,0 +1,105 @@ +/* +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include + +/** + * @addtogroup hipKernelNameRefByPtr hipKernelNameRefByPtr + * @{ + * @ingroup CallbackTest + * `hipKernelNameRefByPtr(const void* hostFunction, hipStream_t stream)` - + * returns the name of passed function pointer on desired stream + */ + +__global__ void test_kernel() { + return; +} + +/** + * Test Description + * ------------------------ + * - Creates new stream and a function pointer + * - Verifies that valid API name is returned + * Test source + * ------------------------ + * - unit/callback/hipKernelNameRefByPtr.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + * - Platform specific (AMD) + */ +TEST_CASE("Unit_hipKernelNameRefByPtr_Positive_Basic") { + const void* kernel_ptr{reinterpret_cast(&test_kernel)}; + + StreamGuard stream_guard{Streams::created}; + REQUIRE(hipKernelNameRefByPtr(kernel_ptr, stream_guard.stream()) != nullptr); +} + +/** + * Test Description + * ------------------------ + * - Passes `nullptr` stream while function pointer is valid + * - Verifies that the returned value is not `nullptr` + * Test source + * ------------------------ + * - unit/callback/hipKernelNameRefByPtr.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + * - Platform specific (AMD) + */ +TEST_CASE("Unit_hipKernelNameRefByPtr_Negative_StreamNullptr") { + const void* kernel_ptr{reinterpret_cast(&test_kernel)}; + StreamGuard stream_guard{Streams::nullstream}; + + REQUIRE(hipKernelNameRefByPtr(kernel_ptr, stream_guard.stream()) != nullptr); +} + +/** + * Test Description + * ------------------------ + * - Performs validation when the function pointer is `nullptr` + * -# When stream is `nullptr` + * - Expected output: return `nullptr` + * -# When stream is valid + * - Expected output: return `nullptr` + * Test source + * ------------------------ + * - unit/callback/hipKernelNameRefByPtr.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + * - Platform specific (AMD) + */ +TEST_CASE("Unit_hipKernelNameRefByPtr_Negative_KernelNullptr") { + const void* kernel_ptr{nullptr}; + + SECTION("stream is nullptr") { + StreamGuard stream_guard{Streams::nullstream}; + REQUIRE(hipKernelNameRefByPtr(kernel_ptr, stream_guard.stream()) == nullptr); + } + + SECTION("stream is created") { + StreamGuard stream_guard{Streams::created}; + REQUIRE(hipKernelNameRefByPtr(kernel_ptr, stream_guard.stream()) == nullptr); + } +}