17 files changed, 340 insertions, 0 deletions
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;
+}