Merge topic 'add_hip_language'

8514ee9b31 HIP: analyze output of `hipcc` to determine default GPU architecture
20d086f1a2 HIP: All HIP tests now run on CMake's current AMD hardware
2e86e50c2f HIP: Add HIP to all the Check* modules
947dbed0aa HIP: Automatically inject the `hip::device` runtime target
b50bfc8913 HIP: Add language to CMake
ff0d2858e1 HIP: Extract clang compiler details from hipcc
bd844387df ROCMClang: Add the ROCm toolkit derived clang compiler to CMake
590553f322 Compilers: protect use of  __has_include
...

Acked-by: Kitware Robot <kwrobot@kitware.com>
Acked-by: Zack Galbreath <zack.galbreath@kitware.com>
Reviewed-by: Raul Tambre <raul@tambre.ee>
Acked-by: Axel Huebl <axel.huebl@plasma.ninja>
Merge-request: !6121

51 files changed, 723 insertions, 5 deletions

diff --git a/Tests/CMakeLists.txt b/Tests/CMakeLists.txt
index fccf15b..57d75e1 100644
--- a/Tests/CMakeLists.txt
+++ b/Tests/CMakeLists.txt
@@ -1485,6 +1485,10 @@ if(BUILD_TESTING)
     add_subdirectory(CudaOnly)
   endif()
 
+  if(CMake_TEST_HIP)
+    add_subdirectory(HIP)
+  endif()
+
   if(CMake_TEST_ISPC)
     add_subdirectory(ISPC)
   endif()
diff --git a/Tests/HIP/ArchitectureOff/CMakeLists.txt b/Tests/HIP/ArchitectureOff/CMakeLists.txt
new file mode 100644
index 0000000..bccb3b4
--- /dev/null
+++ b/Tests/HIP/ArchitectureOff/CMakeLists.txt
@@ -0,0 +1,8 @@
+cmake_minimum_required(VERSION 3.18)
+project(HIPArchitecture HIP)
+
+# Make sure CMake doesn't pass architectures if HIP_ARCHITECTURES is OFF.
+string(APPEND CMAKE_HIP_FLAGS " --offload-arch=gfx908")
+
+add_executable(HIPOnlyArchitectureOff main.hip)
+set_property(TARGET HIPOnlyArchitectureOff PROPERTY HIP_ARCHITECTURES OFF)
diff --git a/Tests/HIP/ArchitectureOff/main.hip b/Tests/HIP/ArchitectureOff/main.hip
new file mode 100644
index 0000000..9256318
--- /dev/null
+++ b/Tests/HIP/ArchitectureOff/main.hip
@@ -0,0 +1,9 @@
+#ifdef __HIP_DEVICE_COMPILE__
+#  ifndef __gfx908__
+#    error "Passed architecture gfx908, but got something else."
+#  endif
+#endif
+
+int main()
+{
+}
diff --git a/Tests/HIP/CMakeLists.txt b/Tests/HIP/CMakeLists.txt
new file mode 100644
index 0000000..9499be8
--- /dev/null
+++ b/Tests/HIP/CMakeLists.txt
@@ -0,0 +1,15 @@
+macro (add_hip_test_macro name)
+  add_test_macro("${name}" ${ARGN})
+  set_property(TEST "${name}" APPEND
+    PROPERTY LABELS "HIP")
+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/CompileFlags/CMakeLists.txt b/Tests/HIP/CompileFlags/CMakeLists.txt
new file mode 100644
index 0000000..c808313
--- /dev/null
+++ b/Tests/HIP/CompileFlags/CMakeLists.txt
@@ -0,0 +1,8 @@
+cmake_minimum_required(VERSION 3.18)
+project(CompileFlags HIP)
+
+add_executable(HIPOnlyCompileFlags main.hip)
+
+set_property(TARGET HIPOnlyCompileFlags PROPERTY HIP_ARCHITECTURES gfx803)
+
+target_compile_options(HIPOnlyCompileFlags PRIVATE -DALWAYS_DEFINE)
diff --git a/Tests/HIP/CompileFlags/main.hip b/Tests/HIP/CompileFlags/main.hip
new file mode 100644
index 0000000..3259f1d
--- /dev/null
+++ b/Tests/HIP/CompileFlags/main.hip
@@ -0,0 +1,13 @@
+#ifdef __HIP_DEVICE_COMPILE__
+#  ifndef __gfx803__
+#    error "Passed architecture gfx803, but got something else."
+#  endif
+#endif
+
+#ifndef ALWAYS_DEFINE
+#  error "ALWAYS_DEFINE not defined!"
+#endif
+
+int main()
+{
+}
diff --git a/Tests/HIP/EnableStandard/CMakeLists.txt b/Tests/HIP/EnableStandard/CMakeLists.txt
new file mode 100644
index 0000000..6701724
--- /dev/null
+++ b/Tests/HIP/EnableStandard/CMakeLists.txt
@@ -0,0 +1,20 @@
+cmake_minimum_required(VERSION 3.18)
+project (EnableStandard HIP)
+
+set(CMAKE_CXX_COMPILER ${CMAKE_HIP_COMPILER})
+enable_language(CXX)
+
+#Goal for this example:
+#build hip sources that require C++11 to be enabled.
+
+add_library(HIPStatic11 STATIC static.cxx)
+set_source_files_properties(static.cxx PROPERTIES LANGUAGE HIP)
+
+add_library(HIPDynamic11 SHARED shared.hip)
+
+add_executable(HIPEnableStandard main.hip)
+target_link_libraries(HIPEnableStandard PRIVATE HIPStatic11 HIPDynamic11)
+
+target_compile_features(HIPDynamic11 PRIVATE cxx_std_11)
+set_target_properties(HIPStatic11 PROPERTIES HIP_STANDARD 11)
+set_target_properties(HIPStatic11 PROPERTIES HIP_STANDARD_REQUIRED TRUE)
diff --git a/Tests/HIP/EnableStandard/main.hip b/Tests/HIP/EnableStandard/main.hip
new file mode 100644
index 0000000..7dc32e8
--- /dev/null
+++ b/Tests/HIP/EnableStandard/main.hip
@@ -0,0 +1,23 @@
+
+#include <iostream>
+
+#ifdef _WIN32
+#  define IMPORT __declspec(dllimport)
+#else
+#  define IMPORT
+#endif
+
+int static_hip11_func(int);
+IMPORT int shared_hip11_func(int);
+
+void test_functions()
+{
+  static_hip11_func(int(42));
+  shared_hip11_func(int(42));
+}
+
+int main(int argc, char** argv)
+{
+  test_functions();
+  return 0;
+}
diff --git a/Tests/HIP/EnableStandard/shared.hip b/Tests/HIP/EnableStandard/shared.hip
new file mode 100644
index 0000000..2042258
--- /dev/null
+++ b/Tests/HIP/EnableStandard/shared.hip
@@ -0,0 +1,15 @@
+
+#include <type_traits>
+
+#ifdef _WIN32
+#  define EXPORT __declspec(dllexport)
+#else
+#  define EXPORT
+#endif
+
+using tt = std::true_type;
+using ft = std::false_type;
+EXPORT int __host__ shared_hip11_func(int x)
+{
+  return x * x + std::integral_constant<int, 17>::value;
+}
diff --git a/Tests/HIP/EnableStandard/static.cxx b/Tests/HIP/EnableStandard/static.cxx
new file mode 100644
index 0000000..2955894
--- /dev/null
+++ b/Tests/HIP/EnableStandard/static.cxx
@@ -0,0 +1,9 @@
+
+#include <type_traits>
+
+using tt = std::true_type;
+using ft = std::false_type;
+int __host__ static_hip11_func(int x)
+{
+  return x * x + std::integral_constant<int, 17>::value;
+}
diff --git a/Tests/HIP/InferHipLang1/CMakeLists.txt b/Tests/HIP/InferHipLang1/CMakeLists.txt
new file mode 100644
index 0000000..63d77fd
--- /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 0000000..6ac8641
--- /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 0000000..987b6c6
--- /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 0000000..0e69de3
--- /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 0000000..6ac8641
--- /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 0000000..987b6c6
--- /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/MathFunctions/CMakeLists.txt b/Tests/HIP/MathFunctions/CMakeLists.txt
new file mode 100644
index 0000000..81e3ddb
--- /dev/null
+++ b/Tests/HIP/MathFunctions/CMakeLists.txt
@@ -0,0 +1,18 @@
+cmake_minimum_required(VERSION 3.18)
+project(MathFunctions HIP)
+
+# This test covers these major HIP language/runtime requirements:
+#
+# 1. This makes sure CMake properly specifies the internal clang header dirs
+#    that hold headers needed for overloads of device side functions
+#
+# 2. This makes sure that all HIP include directories are properly marked as
+#    system includes so we don't get the following warnings:
+#      replacement function 'operator delete' cannot be declared 'inline'#
+#
+# 3. This makes sure CMake properly links to all the built-in libraries
+#    that hip needs that inject support for __half support
+#
+add_executable(HIPOnlyMathFunctions main.hip)
+target_compile_options(HIPOnlyMathFunctions PRIVATE -Werror)
+target_compile_features(HIPOnlyMathFunctions PRIVATE hip_std_14)
diff --git a/Tests/HIP/MathFunctions/main.hip b/Tests/HIP/MathFunctions/main.hip
new file mode 100644
index 0000000..8a6e77f
--- /dev/null
+++ b/Tests/HIP/MathFunctions/main.hip
@@ -0,0 +1,40 @@
+
+#include <stdexcept>
+#include <cmath>
+#include <math.h>
+#include <memory>
+
+#include <hip/hip_runtime.h>
+#include <hip/hip_fp16.h>
+
+namespace {
+template<class T, class F>
+__global__ void global_entry_point(F f, T *out) {
+  *out = f();
+}
+
+template <class T, class F>
+bool verify(F f, T expected)
+{
+  std::unique_ptr<T> cpu_T(new T);
+  T* gpu_T = nullptr;
+  hipMalloc((void**)&gpu_T, sizeof(T));
+  hipLaunchKernelGGL(global_entry_point, 1, 1, 0, 0, f, gpu_T);
+  hipMemcpy(cpu_T.get(), gpu_T, sizeof(T), hipMemcpyDeviceToHost);
+  hipFree(gpu_T);
+  return (*cpu_T == expected);
+}
+}
+
+int main(int argc, char** argv)
+{
+  bool valid = verify([]__device__(){ return std::round(1.4f); }, 1.0f);
+  valid &= verify([]__device__(){ return max<_Float16>(1.0f, 2.0f); }, 2.0f);
+  valid &= verify([]__device__(){ return min<_Float16>(1.0f, 2.0f); }, 1.0f);
+
+  if (valid) {
+    return 0;
+  } else {
+    return 1;
+  }
+}
diff --git a/Tests/HIP/MixedLanguage/CMakeLists.txt b/Tests/HIP/MixedLanguage/CMakeLists.txt
new file mode 100644
index 0000000..4f6dd3b
--- /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 0000000..003d18a
--- /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 0000000..347fba9
--- /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 0000000..8e6c1d3
--- /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 0000000..e6fea9f
--- /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 0000000..06c33b4
--- /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 0000000..2c14fb1
--- /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 0000000..359b9fa
--- /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;
+}
diff --git a/Tests/HIP/TryCompile/CMakeLists.txt b/Tests/HIP/TryCompile/CMakeLists.txt
new file mode 100644
index 0000000..92a834c
--- /dev/null
+++ b/Tests/HIP/TryCompile/CMakeLists.txt
@@ -0,0 +1,15 @@
+cmake_minimum_required(VERSION 3.18)
+project (TryCompile HIP)
+
+#Goal for this example:
+# Verify try_compile with HIP language works
+set(CMAKE_HIP_STANDARD 14)
+set(CMAKE_HIP_ARCHITECTURES gfx803 gfx900)
+
+set(CMAKE_TRY_COMPILE_TARGET_TYPE STATIC_LIBRARY)
+try_compile(result "${CMAKE_CURRENT_BINARY_DIR}"
+        SOURCES "${CMAKE_CURRENT_SOURCE_DIR}/device_function.hip"
+        COPY_FILE "${CMAKE_CURRENT_BINARY_DIR}/device_function.o")
+
+add_executable(HIPOnlyTryCompile main.hip)
+target_link_libraries(HIPOnlyTryCompile "${CMAKE_CURRENT_BINARY_DIR}/device_function.o")
diff --git a/Tests/HIP/TryCompile/device_function.hip b/Tests/HIP/TryCompile/device_function.hip
new file mode 100644
index 0000000..7c1205e
--- /dev/null
+++ b/Tests/HIP/TryCompile/device_function.hip
@@ -0,0 +1,17 @@
+#include <system_error>
+#include <hip/hip_runtime_api.h>
+
+static __global__ void fake_hip_kernel()
+{
+}
+
+int __host__ try_compile_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;
+}
diff --git a/Tests/HIP/TryCompile/main.hip b/Tests/HIP/TryCompile/main.hip
new file mode 100644
index 0000000..091dca3
--- /dev/null
+++ b/Tests/HIP/TryCompile/main.hip
@@ -0,0 +1,8 @@
+int __host__ try_compile_hip_func(int x);
+
+int main(int argc, char** argv)
+{
+  try_compile_hip_func(int(42));
+
+  return 0;
+}
diff --git a/Tests/HIP/WithDefs/CMakeLists.txt b/Tests/HIP/WithDefs/CMakeLists.txt
new file mode 100644
index 0000000..e2db182
--- /dev/null
+++ b/Tests/HIP/WithDefs/CMakeLists.txt
@@ -0,0 +1,37 @@
+
+cmake_minimum_required(VERSION 3.18)
+project (WithDefs HIP)
+
+set(CMAKE_HIP_ARCHITECTURES OFF)
+set(release_compile_defs DEFREL)
+
+#Goal for this example:
+#build a executable that needs to be passed a complex define through add_definitions
+#this verifies we can pass C++ style attributes to hipcc
+add_definitions("-DPACKED_DEFINE=[[gnu::packed]]")
+
+add_executable(HIPOnlyWithDefs main.hip.cpp)
+set_source_files_properties(main.hip.cpp PROPERTIES LANGUAGE HIP)
+
+target_compile_features(HIPOnlyWithDefs PRIVATE hip_std_17)
+
+target_compile_options(HIPOnlyWithDefs
+  PRIVATE
+    --offload-arch=gfx900
+    -DFLAG_COMPILE_LANG_$<COMPILE_LANGUAGE>
+    $<$<HIP_COMPILER_ID:ROCMClang>:-DFLAG_LANG_IS_HIP=$<COMPILE_LANGUAGE:HIP>> # Host-only defines are possible only on NVCC.
+  )
+
+target_compile_definitions(HIPOnlyWithDefs
+  PRIVATE
+    $<$<CONFIG:RELEASE>:$<BUILD_INTERFACE:${release_compile_defs}>>
+    -DDEF_COMPILE_LANG_$<COMPILE_LANGUAGE>
+    -DDEF_LANG_IS_HIP=$<COMPILE_LANGUAGE:HIP>
+    -DDEF_HIP_COMPILER=$<HIP_COMPILER_ID>
+    -DDEF_HIP_COMPILER_VERSION=$<HIP_COMPILER_VERSION>
+  )
+
+target_include_directories(HIPOnlyWithDefs
+  PRIVATE
+    $<$<COMPILE_LANGUAGE:HIP>:${CMAKE_CURRENT_SOURCE_DIR}/inc_hip>
+)
diff --git a/Tests/HIP/WithDefs/inc_hip/inc_hip.h b/Tests/HIP/WithDefs/inc_hip/inc_hip.h
new file mode 100644
index 0000000..fe2698a
--- /dev/null
+++ b/Tests/HIP/WithDefs/inc_hip/inc_hip.h
@@ -0,0 +1 @@
+#define INC_HIP
diff --git a/Tests/HIP/WithDefs/main.hip.cpp b/Tests/HIP/WithDefs/main.hip.cpp
new file mode 100644
index 0000000..a8f2d18
--- /dev/null
+++ b/Tests/HIP/WithDefs/main.hip.cpp
@@ -0,0 +1,89 @@
+#include <iostream>
+
+#include <hip/hip_runtime_api.h>
+#include <inc_hip.h>
+#ifndef INC_HIP
+#  error "INC_HIP not defined!"
+#endif
+
+#ifndef PACKED_DEFINE
+#  error "PACKED_DEFINE not defined!"
+#endif
+
+#ifndef FLAG_COMPILE_LANG_HIP
+#  error "FLAG_COMPILE_LANG_HIP not defined!"
+#endif
+
+#ifndef FLAG_LANG_IS_HIP
+#  error "FLAG_LANG_IS_HIP not defined!"
+#endif
+
+#if !FLAG_LANG_IS_HIP
+#  error "Expected FLAG_LANG_IS_HIP"
+#endif
+
+#ifndef DEF_COMPILE_LANG_HIP
+#  error "DEF_COMPILE_LANG_HIP not defined!"
+#endif
+
+#ifndef DEF_LANG_IS_HIP
+#  error "DEF_LANG_IS_HIP not defined!"
+#endif
+
+#if !DEF_LANG_IS_HIP
+#  error "Expected DEF_LANG_IS_HIP"
+#endif
+
+#ifndef DEF_HIP_COMPILER
+#  error "DEF_HIP_COMPILER not defined!"
+#endif
+
+#ifndef DEF_HIP_COMPILER_VERSION
+#  error "DEF_HIP_COMPILER_VERSION not defined!"
+#endif
+
+static __global__ void DetermineIfValidHIPDevice()
+{
+}
+
+#ifdef _MSC_VER
+#  pragma pack(push, 1)
+#  undef PACKED_DEFINE
+#  define PACKED_DEFINE
+#endif
+struct PACKED_DEFINE result_type
+{
+  bool valid;
+  int value;
+#if defined(NDEBUG) && !defined(DEFREL)
+#  error missing DEFREL flag
+#endif
+};
+#ifdef _MSC_VER
+#  pragma pack(pop)
+#endif
+
+result_type can_launch_kernel()
+{
+  result_type r;
+  DetermineIfValidHIPDevice<<<1, 1>>>();
+  r.valid = (hipSuccess == hipGetLastError());
+  if (r.valid) {
+    r.value = 1;
+  } else {
+    r.value = -1;
+  }
+  return r;
+}
+
+int main(int argc, char** argv)
+{
+  hipError_t err;
+  int nDevices = 0;
+  err = hipGetDeviceCount(&nDevices);
+  if (err != hipSuccess) {
+    std::cerr << hipGetErrorString(err) << std::endl;
+    return 1;
+  }
+  return 0;
+}
diff --git a/Tests/RunCMake/CMakeLists.txt b/Tests/RunCMake/CMakeLists.txt
index 670abbc..27bb929 100644
--- a/Tests/RunCMake/CMakeLists.txt
+++ b/Tests/RunCMake/CMakeLists.txt
@@ -452,6 +452,7 @@ function(add_RunCMake_test_try_compile)
       CMAKE_CXX_STANDARD_DEFAULT
       CMake_TEST_CUDA
       CMake_TEST_ISPC
+      CMake_TEST_HIP
       CMake_TEST_FILESYSTEM_1S
       CMAKE_OBJC_STANDARD_DEFAULT
       CMAKE_OBJCXX_STANDARD_DEFAULT
@@ -603,12 +604,15 @@ add_RunCMake_test(target_include_directories)
 add_RunCMake_test(target_sources)
 add_RunCMake_test(CheckCompilerFlag   -DCMake_TEST_CUDA=${CMake_TEST_CUDA}
                                       -DCMake_TEST_ISPC=${CMake_TEST_ISPC}
-                                      -DCMAKE_Fortran_COMPILER_ID=${CMAKE_Fortran_COMPILER_ID})
+                                      -DCMAKE_Fortran_COMPILER_ID=${CMAKE_Fortran_COMPILER_ID}
+                                      -DCMake_TEST_HIP=${CMake_TEST_HIP})
 add_RunCMake_test(CheckSourceCompiles -DCMake_TEST_CUDA=${CMake_TEST_CUDA}
                                       -DCMake_TEST_ISPC=${CMake_TEST_ISPC}
-                                      -DCMAKE_Fortran_COMPILER_ID=${CMAKE_Fortran_COMPILER_ID})
+                                      -DCMAKE_Fortran_COMPILER_ID=${CMAKE_Fortran_COMPILER_ID}
+                                      -DCMake_TEST_HIP=${CMake_TEST_HIP})
 add_RunCMake_test(CheckSourceRuns     -DCMake_TEST_CUDA=${CMake_TEST_CUDA}
-                                      -DCMAKE_Fortran_COMPILER_ID=${CMAKE_Fortran_COMPILER_ID})
+                                      -DCMAKE_Fortran_COMPILER_ID=${CMAKE_Fortran_COMPILER_ID}
+                                      -DCMake_TEST_HIP=${CMake_TEST_HIP})
 set_property(TEST RunCMake.CheckCompilerFlag
                   RunCMake.CheckSourceCompiles
                   RunCMake.CheckSourceRuns
@@ -616,14 +620,20 @@ set_property(TEST RunCMake.CheckCompilerFlag
 set_property(TEST RunCMake.CheckSourceCompiles
                   RunCMake.CheckCompilerFlag
     APPEND PROPERTY LABELS "ISPC")
+set_property(TEST RunCMake.CheckCompilerFlag
+                  RunCMake.CheckSourceCompiles
+                  RunCMake.CheckSourceRuns
+    APPEND PROPERTY LABELS "HIP")
 add_RunCMake_test(CheckModules)
 add_RunCMake_test(CheckIPOSupported)
 if (CMAKE_SYSTEM_NAME MATCHES "(Linux|Darwin)"
     AND (CMAKE_C_COMPILER_ID MATCHES "Clang|GNU" OR CMAKE_Fortran_COMPILER_ID MATCHES "GNU"))
   add_RunCMake_test(CheckLinkerFlag -DCMAKE_C_COMPILER_ID=${CMAKE_C_COMPILER_ID}
                                     -DCMAKE_Fortran_COMPILER_ID=${CMAKE_Fortran_COMPILER_ID}
-                                    -DCMake_TEST_CUDA=${CMake_TEST_CUDA})
+                                    -DCMake_TEST_CUDA=${CMake_TEST_CUDA}
+                                    -DCMake_TEST_HIP=${CMake_TEST_HIP})
   set_property(TEST RunCMake.CheckLinkerFlag APPEND PROPERTY LABELS "CUDA")
+  set_property(TEST RunCMake.CheckLinkerFlag APPEND PROPERTY LABELS "HIP")
 endif()
 
 
@@ -727,6 +737,9 @@ if("${CMAKE_GENERATOR}" MATCHES "Make|Ninja")
   if(DEFINED CMake_TEST_CUDA)
     list(APPEND CompilerLauncher_ARGS -DCMake_TEST_CUDA=${CMake_TEST_CUDA})
   endif()
+  if(DEFINED CMake_TEST_HIP)
+    list(APPEND CompilerLauncher_ARGS -DCMake_TEST_HIP=${CMake_TEST_HIP})
+  endif()
   if(DEFINED CMake_TEST_ISPC)
     list(APPEND CompilerLauncher_ARGS -DCMake_TEST_ISPC=${CMake_TEST_ISPC})
   endif()
@@ -739,7 +752,7 @@ if("${CMAKE_GENERATOR}" MATCHES "Make|Ninja")
   endif()
   add_RunCMake_test(CompilerLauncher)
   set_property(TEST RunCMake.CompilerLauncher APPEND
-    PROPERTY LABELS "CUDA;ISPC")
+    PROPERTY LABELS "CUDA;HIP;ISPC")
   add_RunCMake_test(ctest_labels_for_subprojects)
   add_RunCMake_test(CompilerArgs)
   add_RunCMake_test(LinkerLauncher)
diff --git a/Tests/RunCMake/CheckCompilerFlag/CheckHIPCompilerFlag.cmake b/Tests/RunCMake/CheckCompilerFlag/CheckHIPCompilerFlag.cmake
new file mode 100644
index 0000000..339ce18
--- /dev/null
+++ b/Tests/RunCMake/CheckCompilerFlag/CheckHIPCompilerFlag.cmake
@@ -0,0 +1,13 @@
+
+enable_language (HIP)
+include(CheckCompilerFlag)
+
+check_compiler_flag(HIP "-_this_is_not_a_flag_" SHOULD_FAIL)
+if(SHOULD_FAIL)
+  message(SEND_ERROR "invalid HIP compile flag didn't fail.")
+endif()
+
+check_compiler_flag(HIP "-DFOO" SHOULD_WORK)
+if(NOT SHOULD_WORK)
+  message(SEND_ERROR "${CMAKE_HIP_COMPILER_ID} compiler flag '-DFOO' check failed")
+endif()
diff --git a/Tests/RunCMake/CheckCompilerFlag/RunCMakeTest.cmake b/Tests/RunCMake/CheckCompilerFlag/RunCMakeTest.cmake
index 7a4e2ce..7ef1860 100644
--- a/Tests/RunCMake/CheckCompilerFlag/RunCMakeTest.cmake
+++ b/Tests/RunCMake/CheckCompilerFlag/RunCMakeTest.cmake
@@ -22,3 +22,7 @@ endif()
 if(CMake_TEST_ISPC)
   run_cmake(CheckISPCCompilerFlag)
 endif()
+
+if(CMake_TEST_HIP)
+  run_cmake(CheckHIPCompilerFlag)
+endif()
diff --git a/Tests/RunCMake/CheckLinkerFlag/CheckHIPLinkerFlag.cmake b/Tests/RunCMake/CheckLinkerFlag/CheckHIPLinkerFlag.cmake
new file mode 100644
index 0000000..3bf3b30
--- /dev/null
+++ b/Tests/RunCMake/CheckLinkerFlag/CheckHIPLinkerFlag.cmake
@@ -0,0 +1,3 @@
+
+set (CHECK_LANGUAGE HIP)
+include ("${CMAKE_CURRENT_SOURCE_DIR}/CheckLinkerFlag.cmake")
diff --git a/Tests/RunCMake/CheckLinkerFlag/RunCMakeTest.cmake b/Tests/RunCMake/CheckLinkerFlag/RunCMakeTest.cmake
index 6ec9148..5e5bff6 100644
--- a/Tests/RunCMake/CheckLinkerFlag/RunCMakeTest.cmake
+++ b/Tests/RunCMake/CheckLinkerFlag/RunCMakeTest.cmake
@@ -16,3 +16,7 @@ endif()
 if (CMake_TEST_CUDA)
   run_cmake(CheckCUDALinkerFlag)
 endif()
+
+if (CMake_TEST_HIP)
+  run_cmake(CheckHIPLinkerFlag)
+endif()
diff --git a/Tests/RunCMake/CheckSourceCompiles/CheckHIPSourceCompiles.cmake b/Tests/RunCMake/CheckSourceCompiles/CheckHIPSourceCompiles.cmake
new file mode 100644
index 0000000..911a0d7
--- /dev/null
+++ b/Tests/RunCMake/CheckSourceCompiles/CheckHIPSourceCompiles.cmake
@@ -0,0 +1,27 @@
+
+enable_language (HIP)
+include(CheckSourceCompiles)
+
+check_source_compiles(HIP "I don't build" SHOULD_FAIL)
+if(SHOULD_FAIL)
+  message(SEND_ERROR "invalid HIP source didn't fail.")
+endif()
+
+check_source_compiles(HIP [=[
+  #include <vector>
+  __device__ int d_func() { }
+  int main() {
+    return 0;
+  }
+]=]
+ SHOULD_BUILD)
+if(NOT SHOULD_BUILD)
+  message(SEND_ERROR "Test fail for valid HIP source.")
+endif()
+
+check_source_compiles(HIP "void l(char const (&x)[2]){}; int main() { l(\"\\n\"); return 0;}"
+ SHOULD_BUILD_COMPLEX)
+
+if(NOT SHOULD_BUILD_COMPLEX)
+  message(SEND_ERROR "Test fail for valid HIP complex source.")
+endif()
diff --git a/Tests/RunCMake/CheckSourceCompiles/RunCMakeTest.cmake b/Tests/RunCMake/CheckSourceCompiles/RunCMakeTest.cmake
index 6e9088f..530f133 100644
--- a/Tests/RunCMake/CheckSourceCompiles/RunCMakeTest.cmake
+++ b/Tests/RunCMake/CheckSourceCompiles/RunCMakeTest.cmake
@@ -23,3 +23,7 @@ endif()
 if(CMake_TEST_ISPC)
   run_cmake(CheckISPCSourceCompiles)
 endif()
+
+if(CMake_TEST_HIP)
+  run_cmake(CheckHIPSourceCompiles)
+endif()
diff --git a/Tests/RunCMake/CheckSourceRuns/CheckHIPSourceRuns.cmake b/Tests/RunCMake/CheckSourceRuns/CheckHIPSourceRuns.cmake
new file mode 100644
index 0000000..d9fb8c2
--- /dev/null
+++ b/Tests/RunCMake/CheckSourceRuns/CheckHIPSourceRuns.cmake
@@ -0,0 +1,21 @@
+
+enable_language (HIP)
+include(CheckSourceRuns)
+
+check_source_runs(HIP "int main() {return 2;}" SHOULD_FAIL)
+if(SHOULD_FAIL)
+  message(SEND_ERROR "HIP check_source_runs succeeded, but should have failed.")
+endif()
+
+check_source_runs(HIP
+[=[
+  #include <vector>
+  __device__ __host__ void fake_function();
+  __host__ int main() {
+    return 0;
+  }
+]=]
+ SHOULD_RUN)
+if(NOT SHOULD_RUN)
+  message(SEND_ERROR "HIP check_source_runs failed for valid HIP executable.")
+endif()
diff --git a/Tests/RunCMake/CheckSourceRuns/RunCMakeTest.cmake b/Tests/RunCMake/CheckSourceRuns/RunCMakeTest.cmake
index c99ac8b..4784103 100644
--- a/Tests/RunCMake/CheckSourceRuns/RunCMakeTest.cmake
+++ b/Tests/RunCMake/CheckSourceRuns/RunCMakeTest.cmake
@@ -19,3 +19,7 @@ endif()
 if (CMake_TEST_CUDA)
   run_cmake(CheckCUDASourceRuns)
 endif()
+
+if (CMake_TEST_HIP)
+  run_cmake(CheckHIPSourceRuns)
+endif()
diff --git a/Tests/RunCMake/CompilerLauncher/HIP-common.cmake b/Tests/RunCMake/CompilerLauncher/HIP-common.cmake
new file mode 100644
index 0000000..53ece78
--- /dev/null
+++ b/Tests/RunCMake/CompilerLauncher/HIP-common.cmake
@@ -0,0 +1,5 @@
+enable_language(HIP)
+enable_language(CXX)
+set(CMAKE_VERBOSE_MAKEFILE TRUE)
+
+add_executable(main main.hip)
diff --git a/Tests/RunCMake/CompilerLauncher/HIP-env-Build-stdout.txt b/Tests/RunCMake/CompilerLauncher/HIP-env-Build-stdout.txt
new file mode 100644
index 0000000..3313e31
--- /dev/null
+++ b/Tests/RunCMake/CompilerLauncher/HIP-env-Build-stdout.txt
@@ -0,0 +1 @@
+.*-E env USED_LAUNCHER=1.*
diff --git a/Tests/RunCMake/CompilerLauncher/HIP-env-launch-Build-stdout.txt b/Tests/RunCMake/CompilerLauncher/HIP-env-launch-Build-stdout.txt
new file mode 100644
index 0000000..3313e31
--- /dev/null
+++ b/Tests/RunCMake/CompilerLauncher/HIP-env-launch-Build-stdout.txt
@@ -0,0 +1 @@
+.*-E env USED_LAUNCHER=1.*
diff --git a/Tests/RunCMake/CompilerLauncher/HIP-env.cmake b/Tests/RunCMake/CompilerLauncher/HIP-env.cmake
new file mode 100644
index 0000000..1bf56ce
--- /dev/null
+++ b/Tests/RunCMake/CompilerLauncher/HIP-env.cmake
@@ -0,0 +1 @@
+include(HIP-common.cmake)
diff --git a/Tests/RunCMake/CompilerLauncher/HIP-launch-Build-stdout.txt b/Tests/RunCMake/CompilerLauncher/HIP-launch-Build-stdout.txt
new file mode 100644
index 0000000..3313e31
--- /dev/null
+++ b/Tests/RunCMake/CompilerLauncher/HIP-launch-Build-stdout.txt
@@ -0,0 +1 @@
+.*-E env USED_LAUNCHER=1.*
diff --git a/Tests/RunCMake/CompilerLauncher/HIP-launch-env.cmake b/Tests/RunCMake/CompilerLauncher/HIP-launch-env.cmake
new file mode 100644
index 0000000..37985a5
--- /dev/null
+++ b/Tests/RunCMake/CompilerLauncher/HIP-launch-env.cmake
@@ -0,0 +1,3 @@
+set(CTEST_USE_LAUNCHERS 1)
+include(CTestUseLaunchers)
+include(HIP-env.cmake)
diff --git a/Tests/RunCMake/CompilerLauncher/HIP-launch.cmake b/Tests/RunCMake/CompilerLauncher/HIP-launch.cmake
new file mode 100644
index 0000000..78fd16b
--- /dev/null
+++ b/Tests/RunCMake/CompilerLauncher/HIP-launch.cmake
@@ -0,0 +1,3 @@
+set(CTEST_USE_LAUNCHERS 1)
+include(CTestUseLaunchers)
+include(HIP.cmake)
diff --git a/Tests/RunCMake/CompilerLauncher/HIP.cmake b/Tests/RunCMake/CompilerLauncher/HIP.cmake
new file mode 100644
index 0000000..9d2577a
--- /dev/null
+++ b/Tests/RunCMake/CompilerLauncher/HIP.cmake
@@ -0,0 +1,2 @@
+set(CMAKE_HIP_COMPILER_LAUNCHER "${CMAKE_COMMAND};-E;env;USED_LAUNCHER=1")
+include(HIP-common.cmake)
diff --git a/Tests/RunCMake/CompilerLauncher/RunCMakeTest.cmake b/Tests/RunCMake/CompilerLauncher/RunCMakeTest.cmake
index 787282a..84d0479 100644
--- a/Tests/RunCMake/CompilerLauncher/RunCMakeTest.cmake
+++ b/Tests/RunCMake/CompilerLauncher/RunCMakeTest.cmake
@@ -29,6 +29,9 @@ endif()
 if(CMake_TEST_Fortran)
   list(APPEND langs Fortran)
 endif()
+if(CMake_TEST_HIP)
+  list(APPEND langs HIP)
+endif()
 if(CMake_TEST_ISPC)
   list(APPEND langs ISPC)
 endif()
diff --git a/Tests/RunCMake/CompilerLauncher/main.hip b/Tests/RunCMake/CompilerLauncher/main.hip
new file mode 100644
index 0000000..f8b643a
--- /dev/null
+++ b/Tests/RunCMake/CompilerLauncher/main.hip
@@ -0,0 +1,4 @@
+int main()
+{
+  return 0;
+}