HIP: Add language to CMake

29 files changed, 372 insertions, 1 deletions

diff --git a/Tests/CMakeLists.txt b/Tests/CMakeLists.txt
index 38b4301..346c2c5 100644
--- a/Tests/CMakeLists.txt
+++ b/Tests/CMakeLists.txt
@@ -1479,6 +1479,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..b3966c9
--- /dev/null
+++ b/Tests/HIP/CMakeLists.txt
@@ -0,0 +1,12 @@
+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.MathFunctions HIPOnlyMathFunctions)
+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/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/TryCompile/CMakeLists.txt b/Tests/HIP/TryCompile/CMakeLists.txt
new file mode 100644
index 0000000..f3bb3bf
--- /dev/null
+++ b/Tests/HIP/TryCompile/CMakeLists.txt
@@ -0,0 +1,14 @@
+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_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..3d4461b
--- /dev/null
+++ b/Tests/HIP/WithDefs/CMakeLists.txt
@@ -0,0 +1,36 @@
+
+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
+    -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 d02d7a2..16487f3 100644
--- a/Tests/RunCMake/CMakeLists.txt
+++ b/Tests/RunCMake/CMakeLists.txt
@@ -724,6 +724,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()
@@ -736,7 +739,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/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;
+}