3 files changed, 127 insertions, 0 deletions
diff --git a/Tests/HIP/WithDefs/CMakeLists.txt b/Tests/HIP/WithDefs/CMakeLists.txt
new file mode 100644
index 000000000..e2db1820a
--- /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 000000000..fe2698a02
--- /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 000000000..a8f2d185b
--- /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;
+}