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.
2023-03-06 04:28:26 +01:00
@@ -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",
@@ -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",
@@ -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)
@@ -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()
@@ -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!");
+}
@@ -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 <hip_test_common.hh>
+
+/**
+ * @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<std::string> 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<uint32_t>::min()), Catch::Equals(kUnknownApi));
+  REQUIRE_THAT(hipApiName(std::numeric_limits<uint32_t>::max()), Catch::Equals(kUnknownApi));
+}
@@ -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 <resource_guards.hh>
+
+/**
+ * @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<std::thread> 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);
+}
@@ -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 <hip_test_common.hh>
+
+/**
+ * @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);
+}
@@ -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 <resource_guards.hh>
+
+/**
+ * @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<const void*>(&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<const void*>(&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);
+  }
+}