EXSWHTEC-114 - Implement tests for hipExtLaunchKernel APIs

catch/hipTestMain/config/config_amd_linux_common.json

@@ -14,6 +14,19 @@
 	   "Unit_hipInit_Negative",
 	   "Unit_hipMemset_Negative_OutOfBoundsPtr",
 	   "Unit_hipDeviceReset_Positive_Basic",
-	   "Unit_hipDeviceReset_Positive_Threaded"
+	   "Unit_hipDeviceReset_Positive_Threaded",
+	   "Unit_hipFuncSetCacheConfig_Positive_Basic",
+	   "Unit_hipFuncSetCacheConfig_Negative_Parameters",
+	   "Unit_hipFuncSetSharedMemConfig_Positive_Basic",
+	   "Unit_hipFuncSetSharedMemConfig_Negative_Parameters",
+	   "Unit_hipFuncSetAttribute_Positive_MaxDynamicSharedMemorySize",
+	   "Unit_hipFuncSetAttribute_Positive_PreferredSharedMemoryCarveout",
+	   "Unit_hipFuncSetAttribute_Positive_Parameters",
+	   "Unit_hipFuncSetAttribute_Negative_Parameters",
+	   "NOTE: The following 4 tests are disabled due to defect - EXSWHTEC-240",
+	   "Unit_hipFuncSetCacheConfig_Negative_Not_Supported",
+	   "Unit_hipFuncSetSharedMemConfig_Negative_Not_Supported",
+	   "Unit_hipFuncSetAttribute_Positive_MaxDynamicSharedMemorySize_Not_Supported",
+	   "Unit_hipFuncSetAttribute_Positive_PreferredSharedMemoryCarveout_Not_Supported"
 	]
 }

catch/hipTestMain/config/config_amd_windows_common.json

@@ -93,6 +93,19 @@
 	   "Note: needs to be enabled when streamPerThread issues are fixed",
 	   "Unit_hipStreamSynchronize_NullStreamAndStreamPerThread",
 	   "Note: intermittent Seg fault failure ",
-	   "Unit_hipGraphAddEventRecordNode_Functional_WithoutFlags"
+	   "Unit_hipGraphAddEventRecordNode_Functional_WithoutFlags",
+	   "Unit_hipFuncSetCacheConfig_Positive_Basic",
+	   "Unit_hipFuncSetCacheConfig_Negative_Parameters",
+	   "Unit_hipFuncSetSharedMemConfig_Positive_Basic",
+	   "Unit_hipFuncSetSharedMemConfig_Negative_Parameters",
+	   "Unit_hipFuncSetAttribute_Positive_MaxDynamicSharedMemorySize",
+	   "Unit_hipFuncSetAttribute_Positive_PreferredSharedMemoryCarveout",
+	   "Unit_hipFuncSetAttribute_Positive_Parameters",
+	   "Unit_hipFuncSetAttribute_Negative_Parameters",
+	   "NOTE: The following 4 tests are disabled due to defect - EXSWHTEC-240",
+	   "Unit_hipFuncSetCacheConfig_Negative_Not_Supported",
+	   "Unit_hipFuncSetSharedMemConfig_Negative_Not_Supported",
+	   "Unit_hipFuncSetAttribute_Positive_MaxDynamicSharedMemorySize_Not_Supported",
+	   "Unit_hipFuncSetAttribute_Positive_PreferredSharedMemoryCarveout_Not_Supported"
 	]
 }

catch/hipTestMain/config/config_nvidia_linux_common.json

@@ -0,0 +1,9 @@
+{
+    "DisabledTests":
+	[
+	   "Unit_hipFuncSetCacheConfig_Negative_Not_Supported",
+	   "Unit_hipFuncSetSharedMemConfig_Negative_Not_Supported",
+	   "Unit_hipFuncSetAttribute_Positive_MaxDynamicSharedMemorySize_Not_Supported",
+	   "Unit_hipFuncSetAttribute_Positive_PreferredSharedMemoryCarveout_Not_Supported"
+	]
+}

catch/hipTestMain/config/config_nvidia_windows_common.json

@@ -0,0 +1,9 @@
+{
+    "DisabledTests":
+	[
+	   "Unit_hipFuncSetCacheConfig_Negative_Not_Supported",
+	   "Unit_hipFuncSetSharedMemConfig_Negative_Not_Supported",
+	   "Unit_hipFuncSetAttribute_Positive_MaxDynamicSharedMemorySize_Not_Supported",
+	   "Unit_hipFuncSetAttribute_Positive_PreferredSharedMemoryCarveout_Not_Supported"
+	]
+}

catch/include/resource_guards.hh

@@ -196,11 +196,11 @@ template <typename T> class DrvArrayAllocGuard {
   const hipExtent extent_;
 };
 
-enum class Streams { nullstream, perThread, created };
+enum class Streams { nullstream, perThread, created, withFlags, withPriority };
 
 class StreamGuard {
  public:
-  StreamGuard(const Streams stream_type) : stream_type_{stream_type} {
+  StreamGuard(const Streams stream_type, unsigned int flags = hipStreamDefault, int priority = 0) : stream_type_{stream_type}, flags_{flags}, priority_{priority} {
     switch (stream_type_) {
       case Streams::nullstream:
         stream_ = nullptr;
@@ -210,6 +210,11 @@ class StreamGuard {
         break;
       case Streams::created:
         HIP_CHECK(hipStreamCreate(&stream_));
+        break;
+      case Streams::withFlags:
+        HIP_CHECK(hipStreamCreateWithFlags(&stream_, flags_));
+      case Streams::withPriority:
+        HIP_CHECK(hipStreamCreateWithPriority(&stream_, flags_, priority_));
     }
   }
 
@@ -226,5 +231,53 @@ class StreamGuard {
 
  private:
   const Streams stream_type_;
+  unsigned int flags_;
+  int priority_;
   hipStream_t stream_;
 };
+
+class EventsGuard {
+public:
+  EventsGuard(size_t N) : events_(N) {
+    for (auto &e : events_) HIP_CHECK(hipEventCreate(&e));
+  }
+
+  EventsGuard(const EventsGuard&) = delete;
+  EventsGuard(EventsGuard&&) = delete;
+
+  ~EventsGuard() {
+    for (auto &e : events_) static_cast<void>(hipEventDestroy(e));
+  }
+
+  hipEvent_t& operator[](int index) { return events_[index]; }
+
+  operator hipEvent_t() const { return events_.at(0); }
+
+  std::vector<hipEvent_t>& event_list() { return events_; }
+
+private:
+  std::vector<hipEvent_t> events_;
+};
+
+class StreamsGuard {
+public:
+  StreamsGuard(size_t N) : streams_(N) {
+    for (auto &s : streams_) HIP_CHECK(hipStreamCreate(&s));
+  }
+
+  StreamsGuard(const StreamsGuard&) = delete;
+  StreamsGuard(StreamsGuard&&) = delete;
+
+  ~StreamsGuard() {
+    for (auto &s : streams_) static_cast<void>(hipStreamDestroy(s));
+  }
+
+  hipStream_t& operator[](int index) { return streams_[index]; }
+
+  operator hipStream_t() const { return streams_.at(0); }
+
+  std::vector<hipStream_t>& stream_list() { return streams_; }
+
+private:
+  std::vector<hipStream_t> streams_;
+};

catch/include/utils.hh

@@ -143,3 +143,9 @@ inline bool DeviceAttributesSupport(const int device, Attributes... attributes)
   };
   return (... && DeviceAttributeSupport(device, attributes));
 }
+
+inline int GetDeviceAttribute(const hipDeviceAttribute_t attr, int device) {
+  int value = 0;
+  HIP_CHECK(hipDeviceGetAttribute(&value, attr, device));
+  return value;
+}

catch/unit/CMakeLists.txt

@@ -37,3 +37,4 @@ add_subdirectory(cooperativeGrps)
 #if(HIP_PLATFORM STREQUAL "amd")
 #add_subdirectory(clock)
 #endif()
+add_subdirectory(executionControl)

catch/unit/executionControl/CMakeLists.txt

@@ -0,0 +1,21 @@
+set(TEST_SRC
+    execution_control_common.cc
+    hipFuncSetCacheConfig.cc
+    hipFuncSetSharedMemConfig.cc
+    hipFuncSetAttribute.cc
+    hipFuncGetAttributes.cc
+    hipLaunchCooperativeKernel.cc
+    hipLaunchCooperativeKernelMultiDevice.cc
+)
+
+if(HIP_PLATFORM MATCHES "amd")
+    set(TEST_SRC ${TEST_SRC}
+        hipExtLaunchKernel.cc
+        hipExtLaunchMultiKernelMultiDevice.cc
+    )
+endif()
+
+hip_add_exe_to_target(NAME ExecutionControlTest
+                      TEST_SRC ${TEST_SRC}
+                      TEST_TARGET_NAME build_tests
+                      COMPILE_OPTIONS -std=c++17)

catch/unit/executionControl/execution_control_common.cc

@@ -0,0 +1,37 @@
+/*
+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 "execution_control_common.hh"
+
+#include <hip_test_common.hh>
+#include <hip/hip_cooperative_groups.h>
+
+__global__ void kernel() {}
+
+__global__ void kernel2() {}
+
+__global__ void kernel_42(int* val) { *val = 42; }
+
+__global__ void coop_kernel() {
+  cooperative_groups::grid_group grid = cooperative_groups::this_grid();
+  grid.sync();
+}

catch/unit/executionControl/execution_control_common.hh

@@ -0,0 +1,31 @@
+/*
+Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#pragma once
+
+__global__ void kernel();
+
+__global__ void kernel2();
+
+__global__ void kernel_42(int* val);
+
+__global__ void coop_kernel();