HIP: Automatically inject the `hip::device` runtime target

Any target that might need to link to hip code needs the `hip::device`
target

15 files changed, 255 insertions, 0 deletions

diff --git a/Tests/HIP/CMakeLists.txt b/Tests/HIP/CMakeLists.txt
index b3966c9ebf..9499be8f6c 100644
--- a/Tests/HIP/CMakeLists.txt
+++ b/Tests/HIP/CMakeLists.txt
@@ -7,6 +7,9 @@ endmacro ()
 add_hip_test_macro(HIP.ArchitectureOff HIPOnlyArchitectureOff)
 add_hip_test_macro(HIP.CompileFlags HIPOnlyCompileFlags)
 add_hip_test_macro(HIP.EnableStandard HIPEnableStandard)
+add_hip_test_macro(HIP.InferHipLang1 HIPInferHipLang1)
+add_hip_test_macro(HIP.InferHipLang2 HIPInferHipLang2)
 add_hip_test_macro(HIP.MathFunctions HIPOnlyMathFunctions)
+add_hip_test_macro(HIP.MixedLanguage HIPMixedLanguage)
 add_hip_test_macro(HIP.TryCompile HIPOnlyTryCompile)
 add_hip_test_macro(HIP.WithDefs HIPOnlyWithDefs)
diff --git a/Tests/HIP/InferHipLang1/CMakeLists.txt b/Tests/HIP/InferHipLang1/CMakeLists.txt
new file mode 100644
index 0000000000..63d77fdaa4
--- /dev/null
+++ b/Tests/HIP/InferHipLang1/CMakeLists.txt
@@ -0,0 +1,12 @@
+cmake_minimum_required(VERSION 3.18)
+project(InferHipLang C CXX HIP)
+
+#Goal for this example:
+#make sure that we understand that HIP is the correct link language
+add_library(InterfaceWithHIP INTERFACE)
+target_sources(InterfaceWithHIP INTERFACE interface.hip main.cxx)
+target_compile_features(InterfaceWithHIP INTERFACE hip_std_14)
+target_compile_features(InterfaceWithHIP INTERFACE cxx_std_11)
+
+add_executable(HIPInferHipLang1 )
+target_link_libraries(HIPInferHipLang1 PRIVATE InterfaceWithHIP)
diff --git a/Tests/HIP/InferHipLang1/interface.hip b/Tests/HIP/InferHipLang1/interface.hip
new file mode 100644
index 0000000000..6ac8641cca
--- /dev/null
+++ b/Tests/HIP/InferHipLang1/interface.hip
@@ -0,0 +1,19 @@
+#include <system_error>
+#include <type_traits>
+#include <hip/hip_runtime_api.h>
+
+static __global__ void fake_hip_kernel()
+{
+}
+
+
+int __host__ interface_hip_func(int x)
+{
+
+  fake_hip_kernel<<<1, 1>>>();
+  bool valid = (hipSuccess == hipGetLastError());
+  if (!valid) {
+    throw std::system_error(ENODEV, std::generic_category(), "no hip device");
+  }
+  return x * x + std::integral_constant<int, 17>::value;
+}
diff --git a/Tests/HIP/InferHipLang1/main.cxx b/Tests/HIP/InferHipLang1/main.cxx
new file mode 100644
index 0000000000..987b6c6516
--- /dev/null
+++ b/Tests/HIP/InferHipLang1/main.cxx
@@ -0,0 +1,19 @@
+
+#include <iostream>
+
+#ifdef __HIP_PLATFORM_HCC__
+#  error "__HIP_PLATFORM_HCC__ propagated to C++ compilation!"
+#endif
+
+#ifdef __HIP_ROCclr__
+#  error "__HIP_ROCclr__ propagated to C++ compilation!"
+#endif
+
+int interface_hip_func(int);
+
+int main(int argc, char** argv)
+{
+  interface_hip_func(int(42));
+
+  return 0;
+}
diff --git a/Tests/HIP/InferHipLang2/CMakeLists.txt b/Tests/HIP/InferHipLang2/CMakeLists.txt
new file mode 100644
index 0000000000..0e69de38fe
--- /dev/null
+++ b/Tests/HIP/InferHipLang2/CMakeLists.txt
@@ -0,0 +1,12 @@
+cmake_minimum_required(VERSION 3.18)
+project(InferHipLang C CXX HIP)
+
+#Goal for this example:
+#make sure that we understand that HIP is the correct link language
+add_library(InterfaceWithHIP OBJECT)
+target_sources(InterfaceWithHIP PRIVATE interface.hip main.cxx)
+target_compile_features(InterfaceWithHIP INTERFACE hip_std_14)
+target_compile_features(InterfaceWithHIP INTERFACE cxx_std_11)
+
+add_executable(HIPInferHipLang2 )
+target_link_libraries(HIPInferHipLang2 PRIVATE InterfaceWithHIP)
diff --git a/Tests/HIP/InferHipLang2/interface.hip b/Tests/HIP/InferHipLang2/interface.hip
new file mode 100644
index 0000000000..6ac8641cca
--- /dev/null
+++ b/Tests/HIP/InferHipLang2/interface.hip
@@ -0,0 +1,19 @@
+#include <system_error>
+#include <type_traits>
+#include <hip/hip_runtime_api.h>
+
+static __global__ void fake_hip_kernel()
+{
+}
+
+
+int __host__ interface_hip_func(int x)
+{
+
+  fake_hip_kernel<<<1, 1>>>();
+  bool valid = (hipSuccess == hipGetLastError());
+  if (!valid) {
+    throw std::system_error(ENODEV, std::generic_category(), "no hip device");
+  }
+  return x * x + std::integral_constant<int, 17>::value;
+}
diff --git a/Tests/HIP/InferHipLang2/main.cxx b/Tests/HIP/InferHipLang2/main.cxx
new file mode 100644
index 0000000000..987b6c6516
--- /dev/null
+++ b/Tests/HIP/InferHipLang2/main.cxx
@@ -0,0 +1,19 @@
+
+#include <iostream>
+
+#ifdef __HIP_PLATFORM_HCC__
+#  error "__HIP_PLATFORM_HCC__ propagated to C++ compilation!"
+#endif
+
+#ifdef __HIP_ROCclr__
+#  error "__HIP_ROCclr__ propagated to C++ compilation!"
+#endif
+
+int interface_hip_func(int);
+
+int main(int argc, char** argv)
+{
+  interface_hip_func(int(42));
+
+  return 0;
+}
diff --git a/Tests/HIP/MixedLanguage/CMakeLists.txt b/Tests/HIP/MixedLanguage/CMakeLists.txt
new file mode 100644
index 0000000000..4f6dd3b0fb
--- /dev/null
+++ b/Tests/HIP/MixedLanguage/CMakeLists.txt
@@ -0,0 +1,19 @@
+cmake_minimum_required(VERSION 3.18)
+project (MixedLanguage C CXX HIP)
+
+set(CMAKE_HIP_STANDARD 14)
+set(CMAKE_CXX_STANDARD 14)
+
+#Goal for this example:
+#make sure that we can build multiple languages into targets
+#and have the link language always be HIP
+add_library(MixedSharedLib SHARED shared.c)
+add_library(MixedObjectLib OBJECT shared.cxx shared.hip)
+set_target_properties(MixedObjectLib PROPERTIES POSITION_INDEPENDENT_CODE ON)
+target_link_libraries(MixedSharedLib PRIVATE MixedObjectLib)
+
+add_library(MixedStaticLib STATIC static.c static.cxx static.hip)
+set_target_properties(MixedStaticLib PROPERTIES POSITION_INDEPENDENT_CODE ON)
+
+add_executable(HIPMixedLanguage main.cxx)
+target_link_libraries(HIPMixedLanguage PRIVATE MixedStaticLib MixedSharedLib)
diff --git a/Tests/HIP/MixedLanguage/main.cxx b/Tests/HIP/MixedLanguage/main.cxx
new file mode 100644
index 0000000000..003d18ac4c
--- /dev/null
+++ b/Tests/HIP/MixedLanguage/main.cxx
@@ -0,0 +1,40 @@
+
+#include <iostream>
+
+#ifdef __HIP_PLATFORM_HCC__
+#  error "__HIP_PLATFORM_HCC__ propagated to C++ compilation!"
+#endif
+
+#ifdef __HIP_ROCclr__
+#  error "__HIP_ROCclr__ propagated to C++ compilation!"
+#endif
+
+#ifdef _WIN32
+#  define IMPORT __declspec(dllimport)
+#else
+#  define IMPORT
+#endif
+
+extern "C" {
+IMPORT int shared_c_func(int);
+int static_c_func(int);
+}
+
+IMPORT int shared_cxx_func(int);
+IMPORT int shared_hip_func(int);
+
+int static_cxx_func(int);
+int static_hip_func(int);
+
+int main(int argc, char** argv)
+{
+  static_c_func(int(42));
+  static_cxx_func(int(42));
+  static_hip_func(int(42));
+
+  shared_c_func(int(42));
+  shared_cxx_func(int(42));
+  shared_hip_func(int(42));
+
+  return 0;
+}
diff --git a/Tests/HIP/MixedLanguage/shared.c b/Tests/HIP/MixedLanguage/shared.c
new file mode 100644
index 0000000000..347fba9255
--- /dev/null
+++ b/Tests/HIP/MixedLanguage/shared.c
@@ -0,0 +1,12 @@
+
+
+#ifdef _WIN32
+#  define EXPORT __declspec(dllexport)
+#else
+#  define EXPORT
+#endif
+
+EXPORT int shared_c_func(int x)
+{
+  return -x;
+}
diff --git a/Tests/HIP/MixedLanguage/shared.cxx b/Tests/HIP/MixedLanguage/shared.cxx
new file mode 100644
index 0000000000..8e6c1d35f8
--- /dev/null
+++ b/Tests/HIP/MixedLanguage/shared.cxx
@@ -0,0 +1,21 @@
+
+#include <type_traits>
+
+#ifdef __HIP_PLATFORM_HCC__
+#  error "__HIP_PLATFORM_HCC__ propagated to C++ compilation!"
+#endif
+
+#ifdef __HIP_ROCclr__
+#  error "__HIP_ROCclr__ propagated to C++ compilation!"
+#endif
+
+#ifdef _WIN32
+#  define EXPORT __declspec(dllexport)
+#else
+#  define EXPORT
+#endif
+
+EXPORT int shared_cxx_func(int x)
+{
+  return x * x + std::integral_constant<int, 14>::value;
+}
diff --git a/Tests/HIP/MixedLanguage/shared.hip b/Tests/HIP/MixedLanguage/shared.hip
new file mode 100644
index 0000000000..e6fea9fe21
--- /dev/null
+++ b/Tests/HIP/MixedLanguage/shared.hip
@@ -0,0 +1,26 @@
+#include <system_error>
+#include <type_traits>
+#include <hip/hip_runtime_api.h>
+
+#ifdef _WIN32
+#  define EXPORT __declspec(dllexport)
+#else
+#  define EXPORT
+#endif
+
+
+static __global__ void fake_hip_kernel()
+{
+}
+
+
+int __host__ shared_hip_func(int x)
+{
+
+  fake_hip_kernel<<<1, 1>>>();
+  bool valid = (hipSuccess == hipGetLastError());
+  if (!valid) {
+    throw std::system_error(ENODEV, std::generic_category(), "no hip device");
+  }
+  return x * x + std::integral_constant<int, 17>::value;
+}
diff --git a/Tests/HIP/MixedLanguage/static.c b/Tests/HIP/MixedLanguage/static.c
new file mode 100644
index 0000000000..06c33b4cc7
--- /dev/null
+++ b/Tests/HIP/MixedLanguage/static.c
@@ -0,0 +1,6 @@
+
+
+int static_c_func(int x)
+{
+  return -x;
+}
diff --git a/Tests/HIP/MixedLanguage/static.cxx b/Tests/HIP/MixedLanguage/static.cxx
new file mode 100644
index 0000000000..2c14fb1018
--- /dev/null
+++ b/Tests/HIP/MixedLanguage/static.cxx
@@ -0,0 +1,7 @@
+
+#include <type_traits>
+
+int static_cxx_func(int x)
+{
+  return x * x + std::integral_constant<int, 14>::value;
+}
diff --git a/Tests/HIP/MixedLanguage/static.hip b/Tests/HIP/MixedLanguage/static.hip
new file mode 100644
index 0000000000..359b9facda
--- /dev/null
+++ b/Tests/HIP/MixedLanguage/static.hip
@@ -0,0 +1,21 @@
+
+#include <type_traits>
+#include <system_error>
+#include <hip/hip_runtime_api.h>
+
+
+static __global__ void fake_hip_kernel()
+{
+}
+
+
+int __host__ static_hip_func(int x)
+{
+
+  fake_hip_kernel<<<1, 1>>>();
+  bool valid = (hipSuccess == hipGetLastError());
+  if (!valid) {
+    throw std::system_error(ENODEV, std::generic_category(), "no hip device");
+  }
+  return x * x + std::integral_constant<int, 17>::value;
+}