HIP: Add language to CMake

Help/command/enable_language.rst

@@ -10,7 +10,7 @@ Enables support for the named language in CMake.  This is
 the same as the :command:`project` command but does not create any of the extra
 variables that are created by the project command.  Example languages
 are ``CXX``, ``C``, ``CUDA``, ``OBJC``, ``OBJCXX``, ``Fortran``,
-``ISPC``, and ``ASM``.
+``HIP``, ``ISPC``, and ``ASM``.
 
 .. versionadded:: 3.8
   Added ``CUDA`` support.
@@ -21,6 +21,9 @@ are ``CXX``, ``C``, ``CUDA``, ``OBJC``, ``OBJCXX``, ``Fortran``,
 .. versionadded:: 3.18
   Added ``ISPC`` support.
 
+.. versionadded:: 3.21
+  Added ``HIP`` support.
+
 If enabling ``ASM``, enable it last so that CMake can check whether
 compilers for other languages like ``C`` work for assembly too.
 

Help/command/project.rst

@@ -103,7 +103,7 @@ The options are:
 
   Selects which programming languages are needed to build the project.
   Supported languages include ``C``, ``CXX`` (i.e.  C++), ``CUDA``,
-  ``OBJC`` (i.e. Objective-C), ``OBJCXX``, ``Fortran``, ``ISPC``, and ``ASM``.
+  ``OBJC`` (i.e. Objective-C), ``OBJCXX``, ``Fortran``, ``HIP``, ``ISPC``, and ``ASM``.
   By default ``C`` and ``CXX`` are enabled if no language options are given.
   Specify language ``NONE``, or use the ``LANGUAGES`` keyword and list no languages,
   to skip enabling any languages.

Help/envvar/HIPCXX.rst

@@ -0,0 +1,13 @@
+HIPCXX
+------
+
+.. versionadded:: 3.21
+
+.. include:: ENV_VAR.txt
+
+Preferred executable for compiling ``HIP`` language files. Will only be used by
+CMake on the first configuration to determine ``HIP`` compiler, after which the
+value for ``HIP`` is stored in the cache as
+:variable:`CMAKE_HIP_COMPILER <CMAKE_<LANG>_COMPILER>`. For any configuration
+run (including the first), the environment variable will be ignored if the
+:variable:`CMAKE_HIP_COMPILER <CMAKE_<LANG>_COMPILER>` variable is defined.

Help/envvar/HIPFLAGS.rst

@@ -0,0 +1,15 @@
+HIPFLAGS
+--------
+
+.. versionadded:: 3.21
+
+.. include:: ENV_VAR.txt
+
+Default compilation flags to be used when compiling ``HIP`` files. Will only be
+used by CMake on the first configuration to determine ``HIP`` default
+compilation flags, after which the value for ``HIPFLAGS`` is stored in the
+cache as :variable:`CMAKE_HIP_FLAGS <CMAKE_<LANG>_FLAGS>`. For any configuration
+run (including the first), the environment variable will be ignored if
+the :variable:`CMAKE_HIP_FLAGS <CMAKE_<LANG>_FLAGS>` variable is defined.
+
+See also :variable:`CMAKE_HIP_FLAGS_INIT <CMAKE_<LANG>_FLAGS_INIT>`.

Help/manual/cmake-env-variables.7.rst

@@ -67,6 +67,8 @@ Environment Variables for Languages
    /envvar/CXXFLAGS
    /envvar/FC
    /envvar/FFLAGS
+   /envvar/HIPCXX
+   /envvar/HIPFLAGS
    /envvar/ISPC
    /envvar/ISPCFLAGS
    /envvar/OBJC

Help/manual/cmake-generator-expressions.7.rst

@@ -181,6 +181,13 @@ Variable Queries
   of the entries in ``compiler_ids``, otherwise ``0``.
   See also the :variable:`CMAKE_<LANG>_COMPILER_ID` variable.
 
+.. genex:: $<HIP_COMPILER_ID:compiler_ids>
+
+  where ``compiler_ids`` is a comma-separated list.
+  ``1`` if the CMake's compiler id of the HIP compiler matches any one
+  of the entries in ``compiler_ids``, otherwise ``0``.
+  See also the :variable:`CMAKE_<LANG>_COMPILER_ID` variable.
+
 .. genex:: $<ISPC_COMPILER_ID:compiler_ids>
 
   where ``compiler_ids`` is a comma-separated list.
@@ -218,6 +225,11 @@ Variable Queries
   ``1`` if the version of the Fortran compiler matches ``version``, otherwise ``0``.
   See also the :variable:`CMAKE_<LANG>_COMPILER_VERSION` variable.
 
+.. genex:: $<HIP_COMPILER_VERSION:version>
+
+  ``1`` if the version of the HIP compiler matches ``version``, otherwise ``0``.
+  See also the :variable:`CMAKE_<LANG>_COMPILER_VERSION` variable.
+
 .. genex:: $<ISPC_COMPILER_VERSION:version>
 
   ``1`` if the version of the ISPC compiler matches ``version``, otherwise ``0``.
@@ -648,6 +660,11 @@ Variable Queries
   The CMake's compiler id of the Fortran compiler used.
   See also the :variable:`CMAKE_<LANG>_COMPILER_ID` variable.
 
+.. genex:: $<HIP_COMPILER_ID>
+
+  The CMake's compiler id of the HIP compiler used.
+  See also the :variable:`CMAKE_<LANG>_COMPILER_ID` variable.
+
 .. genex:: $<ISPC_COMPILER_ID>
 
   The CMake's compiler id of the ISPC compiler used.
@@ -683,6 +700,11 @@ Variable Queries
   The version of the Fortran compiler used.
   See also the :variable:`CMAKE_<LANG>_COMPILER_VERSION` variable.
 
+.. genex:: $<HIP_COMPILER_VERSION>
+
+  The version of the HIP compiler used.
+  See also the :variable:`CMAKE_<LANG>_COMPILER_VERSION` variable.
+
 .. genex:: $<ISPC_COMPILER_VERSION>
 
   The version of the ISPC compiler used.

Help/manual/cmake-properties.7.rst

@@ -212,6 +212,7 @@ Properties on Targets
    /prop_tgt/GHS_NO_SOURCE_GROUP_FILE
    /prop_tgt/GNUtoMS
    /prop_tgt/HAS_CXX
+   /prop_tgt/HIP_ARCHITECTURES
    /prop_tgt/IMPLICIT_DEPENDS_INCLUDE_TRANSFORM
    /prop_tgt/IMPORTED
    /prop_tgt/IMPORTED_COMMON_LANGUAGE_RUNTIME

Help/manual/cmake-variables.7.rst

@@ -523,6 +523,7 @@ Variables for Languages
    /variable/CMAKE_Fortran_MODDIR_DEFAULT
    /variable/CMAKE_Fortran_MODDIR_FLAG
    /variable/CMAKE_Fortran_MODOUT_FLAG
+   /variable/CMAKE_HIP_ARCHITECTURES
    /variable/CMAKE_ISPC_HEADER_DIRECTORY
    /variable/CMAKE_ISPC_HEADER_SUFFIX
    /variable/CMAKE_ISPC_INSTRUCTION_SETS

Help/prop_sf/LANGUAGE.rst

@@ -6,8 +6,8 @@ Specify the programming language in which a source file is written.
 A property that can be set to indicate what programming language the
 source file is.  If it is not set the language is determined based on
 the file extension.  Typical values are ``CXX`` (i.e.  C++), ``C``,
-``CSharp``, ``CUDA``, ``Fortran``, ``ISPC``, and ``ASM``.  Setting this
-property for a file means this file will be compiled.  Do not set this
+``CSharp``, ``CUDA``, ``Fortran``, ``HIP``, ``ISPC``, and ``ASM``.  Setting
+this property for a file means this file will be compiled.  Do not set this
 for headers or files that should not be compiled.
 
 .. versionchanged:: 3.20

Help/prop_tgt/HIP_ARCHITECTURES.rst

@@ -0,0 +1,26 @@
+HIP_ARCHITECTURES
+-----------------
+
+.. versionadded:: 3.21
+
+List of AMD GPU architectures to generate device code for.
+
+An empty or false value (e.g. ``OFF``) defers architecture generation to compiler
+defaults.
+
+This property is initialized by the value of the :variable:`CMAKE_HIP_ARCHITECTURES`
+variable if it is set when a target is created.
+
+The HIP compilation model has two modes: whole and separable. Whole compilation
+generates device code at compile time. Separable compilation generates device
+code at link time. Therefore the ``HIP_ARCHITECTURES`` target property should
+be set on targets that compile or link with any HIP sources.
+
+Examples
+^^^^^^^^
+
+.. code-block:: cmake
+
+  set_property(TARGET tgt PROPERTY HIP_ARCHITECTURES gfx801 gfx900)
+
+Generates code for both ``gfx801`` and ``gfx900``.

Help/prop_tgt/LANG_COMPILER_LAUNCHER.rst

@@ -4,7 +4,7 @@
 .. versionadded:: 3.4
 
 This property is implemented only when ``<LANG>`` is ``C``, ``CXX``,
-``Fortran``, ``ISPC``, ``OBJC``, ``OBJCXX``, or ``CUDA``.
+``Fortran``, ``HIP``, ``ISPC``, ``OBJC``, ``OBJCXX``, or ``CUDA``.
 
 Specify a :ref:`semicolon-separated list <CMake Language Lists>` containing a command line
 for a compiler launching tool. The :ref:`Makefile Generators` and the

Help/release/dev/hip.rst

@@ -0,0 +1,5 @@
+hip
+---
+
+* CMake learned to support ``HIP`` as a first-class language that can be
+  enabled via the :command:`project` and :command:`enable_language` commands.

Help/variable/CMAKE_HIP_ARCHITECTURES.rst

@@ -0,0 +1,11 @@
+CMAKE_HIP_ARCHITECTURES
+-----------------------
+
+.. versionadded:: 3.21
+
+Default value for :prop_tgt:`HIP_ARCHITECTURES` property of targets.
+
+This is initialized to ``OFF``.
+
+This variable is used to initialize the :prop_tgt:`HIP_ARCHITECTURES` property
+on all targets. See the target property for additional information.

Help/variable/CMAKE_LANG_COMPILER_LAUNCHER.rst

@@ -6,7 +6,7 @@ CMAKE_<LANG>_COMPILER_LAUNCHER
 Default value for :prop_tgt:`<LANG>_COMPILER_LAUNCHER` target property.
 This variable is used to initialize the property on each target as it is
 created.  This is done only when ``<LANG>`` is ``C``, ``CXX``, ``Fortran``,
-``ISPC``, ``OBJC``, ``OBJCXX``, or ``CUDA``.
+``HIP``, ``ISPC``, ``OBJC``, ``OBJCXX``, or ``CUDA``.
 
 This variable is initialized to the :envvar:`CMAKE_<LANG>_COMPILER_LAUNCHER`
 environment variable if it is set.

Modules/CMakeDetermineCompileFeatures.cmake

@@ -166,6 +166,62 @@ function(cmake_determine_compile_features lang)
 
     message(CHECK_PASS "done")
 
+  elseif(lang STREQUAL HIP AND COMMAND cmake_record_hip_compile_features)
+    message(CHECK_START "Detecting ${lang} compile features")
+
+    set(CMAKE_HIP98_COMPILE_FEATURES)
+    set(CMAKE_HIP11_COMPILE_FEATURES)
+    set(CMAKE_HIP14_COMPILE_FEATURES)
+    set(CMAKE_HIP17_COMPILE_FEATURES)
+    set(CMAKE_HIP20_COMPILE_FEATURES)
+    set(CMAKE_HIP23_COMPILE_FEATURES)
+
+    include("${CMAKE_ROOT}/Modules/Internal/FeatureTesting.cmake")
+
+    cmake_record_hip_compile_features()
+
+    if(NOT _result EQUAL 0)
+      message(CHECK_FAIL "failed")
+      return()
+    endif()
+
+    if (CMAKE_HIP20_COMPILE_FEATURES AND CMAKE_HIP23_COMPILE_FEATURES)
+      list(REMOVE_ITEM CMAKE_HIP23_COMPILE_FEATURES ${CMAKE_HIP20_COMPILE_FEATURES})
+    endif()
+    if (CMAKE_HIP17_COMPILE_FEATURES AND CMAKE_HIP20_COMPILE_FEATURES)
+      list(REMOVE_ITEM CMAKE_HIP20_COMPILE_FEATURES ${CMAKE_HIP17_COMPILE_FEATURES})
+    endif()
+    if (CMAKE_HIP14_COMPILE_FEATURES AND CMAKE_HIP17_COMPILE_FEATURES)
+      list(REMOVE_ITEM CMAKE_HIP17_COMPILE_FEATURES ${CMAKE_HIP14_COMPILE_FEATURES})
+    endif()
+    if (CMAKE_HIP11_COMPILE_FEATURES AND CMAKE_HIP14_COMPILE_FEATURES)
+      list(REMOVE_ITEM CMAKE_HIP14_COMPILE_FEATURES ${CMAKE_HIP11_COMPILE_FEATURES})
+    endif()
+    if (CMAKE_HIP98_COMPILE_FEATURES AND CMAKE_HIP11_COMPILE_FEATURES)
+      list(REMOVE_ITEM CMAKE_HIP11_COMPILE_FEATURES ${CMAKE_HIP98_COMPILE_FEATURES})
+    endif()
+
+    if(NOT CMAKE_HIP_COMPILE_FEATURES)
+      set(CMAKE_HIP_COMPILE_FEATURES
+        ${CMAKE_HIP98_COMPILE_FEATURES}
+        ${CMAKE_HIP11_COMPILE_FEATURES}
+        ${CMAKE_HIP14_COMPILE_FEATURES}
+        ${CMAKE_HIP17_COMPILE_FEATURES}
+        ${CMAKE_HIP20_COMPILE_FEATURES}
+        ${CMAKE_HIP23_COMPILE_FEATURES}
+      )
+    endif()
+
+    set(CMAKE_HIP_COMPILE_FEATURES ${CMAKE_HIP_COMPILE_FEATURES} PARENT_SCOPE)
+    set(CMAKE_HIP98_COMPILE_FEATURES ${CMAKE_HIP98_COMPILE_FEATURES} PARENT_SCOPE)
+    set(CMAKE_HIP11_COMPILE_FEATURES ${CMAKE_HIP11_COMPILE_FEATURES} PARENT_SCOPE)
+    set(CMAKE_HIP14_COMPILE_FEATURES ${CMAKE_HIP14_COMPILE_FEATURES} PARENT_SCOPE)
+    set(CMAKE_HIP17_COMPILE_FEATURES ${CMAKE_HIP17_COMPILE_FEATURES} PARENT_SCOPE)
+    set(CMAKE_HIP20_COMPILE_FEATURES ${CMAKE_HIP20_COMPILE_FEATURES} PARENT_SCOPE)
+    set(CMAKE_HIP23_COMPILE_FEATURES ${CMAKE_HIP23_COMPILE_FEATURES} PARENT_SCOPE)
+
+    message(CHECK_PASS "done")
+
   endif()
 
 endfunction()

Modules/CMakeDetermineHIPCompiler.cmake

@@ -0,0 +1,96 @@
+# Distributed under the OSI-approved BSD 3-Clause License.  See accompanying
+# file Copyright.txt or https://cmake.org/licensing for details.
+
+include(${CMAKE_ROOT}/Modules/CMakeDetermineCompiler.cmake)
+include(${CMAKE_ROOT}/Modules/CMakeParseImplicitLinkInfo.cmake)
+
+if( NOT ( ("${CMAKE_GENERATOR}" MATCHES "Make") OR
+          ("${CMAKE_GENERATOR}" MATCHES "Ninja") ) )
+  message(FATAL_ERROR "HIP language not currently supported by \"${CMAKE_GENERATOR}\" generator")
+endif()
+
+
+if(NOT CMAKE_HIP_COMPILER)
+  set(CMAKE_HIP_COMPILER_INIT NOTFOUND)
+
+  # prefer the environment variable HIPCXX
+  if(NOT $ENV{HIPCXX} STREQUAL "")
+    get_filename_component(CMAKE_HIP_COMPILER_INIT $ENV{HIPCXX} PROGRAM PROGRAM_ARGS CMAKE_HIP_FLAGS_ENV_INIT)
+    if(CMAKE_HIP_FLAGS_ENV_INIT)
+      set(CMAKE_HIP_COMPILER_ARG1 "${CMAKE_HIP_FLAGS_ENV_INIT}" CACHE STRING "Arguments to CXX compiler")
+    endif()
+    if(NOT EXISTS ${CMAKE_HIP_COMPILER_INIT})
+      message(FATAL_ERROR "Could not find compiler set in environment variable HIPCXX:\n$ENV{HIPCXX}.\n${CMAKE_HIP_COMPILER_INIT}")
+    endif()
+  endif()
+
+  # finally list compilers to try
+  if(NOT CMAKE_HIP_COMPILER_INIT)
+    set(CMAKE_HIP_COMPILER_LIST hipcc clang++)
+  endif()
+
+  _cmake_find_compiler(HIP)
+else()
+  _cmake_find_compiler_path(HIP)
+endif()
+
+mark_as_advanced(CMAKE_HIP_COMPILER)
+
+# Build a small source file to identify the compiler.
+if(NOT CMAKE_HIP_COMPILER_ID_RUN)
+  set(CMAKE_HIP_COMPILER_ID_RUN 1)
+
+  # Try to identify the compiler.
+  set(CMAKE_HIP_COMPILER_ID)
+  set(CMAKE_HIP_PLATFORM_ID)
+  file(READ ${CMAKE_ROOT}/Modules/CMakePlatformId.h.in
+    CMAKE_HIP_COMPILER_ID_PLATFORM_CONTENT)
+
+  list(APPEND CMAKE_HIP_COMPILER_ID_TEST_FLAGS_FIRST "-v")
+
+  include(${CMAKE_ROOT}/Modules/CMakeDetermineCompilerId.cmake)
+  CMAKE_DETERMINE_COMPILER_ID(HIP HIPFLAGS CMakeHIPCompilerId.hip)
+
+  _cmake_find_compiler_sysroot(HIP)
+
+endif()
+
+if (NOT _CMAKE_TOOLCHAIN_LOCATION)
+  get_filename_component(_CMAKE_TOOLCHAIN_LOCATION "${CMAKE_HIP_COMPILER}" PATH)
+endif ()
+
+set(_CMAKE_PROCESSING_LANGUAGE "HIP")
+include(CMakeFindBinUtils)
+include(Compiler/${CMAKE_HIP_COMPILER_ID}-FindBinUtils OPTIONAL)
+unset(_CMAKE_PROCESSING_LANGUAGE)
+
+if(CMAKE_HIP_COMPILER_SYSROOT)
+  string(CONCAT _SET_CMAKE_HIP_COMPILER_SYSROOT
+    "set(CMAKE_HIP_COMPILER_SYSROOT \"${CMAKE_HIP_COMPILER_SYSROOT}\")\n"
+    "set(CMAKE_COMPILER_SYSROOT \"${CMAKE_HIP_COMPILER_SYSROOT}\")")
+else()
+  set(_SET_CMAKE_HIP_COMPILER_SYSROOT "")
+endif()
+
+if(CMAKE_HIP_COMPILER_ARCHITECTURE_ID)
+  set(_SET_CMAKE_HIP_COMPILER_ARCHITECTURE_ID
+    "set(CMAKE_HIP_COMPILER_ARCHITECTURE_ID ${CMAKE_HIP_COMPILER_ARCHITECTURE_ID})")
+else()
+  set(_SET_CMAKE_HIP_COMPILER_ARCHITECTURE_ID "")
+endif()
+
+if(MSVC_HIP_ARCHITECTURE_ID)
+  set(SET_MSVC_HIP_ARCHITECTURE_ID
+    "set(MSVC_HIP_ARCHITECTURE_ID ${MSVC_HIP_ARCHITECTURE_ID})")
+endif()
+
+if(NOT DEFINED CMAKE_HIP_ARCHITECTURES)
+  set(CMAKE_HIP_ARCHITECTURES "OFF" CACHE STRING "HIP architectures")
+endif()
+
+# configure variables set in this file for fast reload later on
+configure_file(${CMAKE_ROOT}/Modules/CMakeHIPCompiler.cmake.in
+  ${CMAKE_PLATFORM_INFO_DIR}/CMakeHIPCompiler.cmake
+  @ONLY
+  )
+set(CMAKE_HIP_COMPILER_ENV_VAR "HIPCXX")

Modules/CMakeHIPCompiler.cmake.in

@@ -0,0 +1,59 @@
+set(CMAKE_HIP_COMPILER "@CMAKE_HIP_COMPILER@")
+set(CMAKE_HIP_COMPILER_ID "@CMAKE_HIP_COMPILER_ID@")
+set(CMAKE_HIP_COMPILER_VERSION "@CMAKE_HIP_COMPILER_VERSION@")
+set(CMAKE_HIP_STANDARD_COMPUTED_DEFAULT "@CMAKE_HIP_STANDARD_COMPUTED_DEFAULT@")
+set(CMAKE_HIP_COMPILE_FEATURES "@CMAKE_HIP_COMPILE_FEATURES@")
+set(CMAKE_HIP98_COMPILE_FEATURES "@CMAKE_HIP03_COMPILE_FEATURES@")
+set(CMAKE_HIP11_COMPILE_FEATURES "@CMAKE_HIP11_COMPILE_FEATURES@")
+set(CMAKE_HIP14_COMPILE_FEATURES "@CMAKE_HIP14_COMPILE_FEATURES@")
+set(CMAKE_HIP17_COMPILE_FEATURES "@CMAKE_HIP17_COMPILE_FEATURES@")
+set(CMAKE_HIP20_COMPILE_FEATURES "@CMAKE_HIP20_COMPILE_FEATURES@")
+set(CMAKE_HIP23_COMPILE_FEATURES "@CMAKE_HIP23_COMPILE_FEATURES@")
+
+set(CMAKE_HIP_PLATFORM_ID "@CMAKE_HIP_PLATFORM_ID@")
+set(CMAKE_HIP_SIMULATE_ID "@CMAKE_HIP_SIMULATE_ID@")
+set(CMAKE_HIP_COMPILER_FRONTEND_VARIANT "@CMAKE_HIP_COMPILER_FRONTEND_VARIANT@")
+set(CMAKE_HIP_SIMULATE_VERSION "@CMAKE_HIP_SIMULATE_VERSION@")
+@SET_MSVC_HIP_ARCHITECTURE_ID@
+@_SET_CMAKE_HIP_COMPILER_SYSROOT@
+
+set(CMAKE_HIP_COMPILER_ENV_VAR "HIPCXX")
+
+set(CMAKE_HIP_COMPILER_LOADED 1)
+set(CMAKE_HIP_COMPILER_ID_RUN 1)
+set(CMAKE_HIP_SOURCE_FILE_EXTENSIONS hip)
+set(CMAKE_HIP_LINKER_PREFERENCE 15)
+set(CMAKE_HIP_LINKER_PREFERENCE_PROPAGATES 1)
+
+set(CMAKE_HIP_SIZEOF_DATA_PTR "@CMAKE_HIP_SIZEOF_DATA_PTR@")
+set(CMAKE_HIP_COMPILER_ABI "@CMAKE_HIP_COMPILER_ABI@")
+set(CMAKE_HIP_LIBRARY_ARCHITECTURE "@CMAKE_HIP_LIBRARY_ARCHITECTURE@")
+
+if(CMAKE_HIP_SIZEOF_DATA_PTR)
+  set(CMAKE_SIZEOF_VOID_P "${CMAKE_HIP_SIZEOF_DATA_PTR}")
+endif()
+
+if(CMAKE_HIP_COMPILER_ABI)
+  set(CMAKE_INTERNAL_PLATFORM_ABI "${CMAKE_HIP_COMPILER_ABI}")
+endif()
+
+if(CMAKE_HIP_LIBRARY_ARCHITECTURE)
+  set(CMAKE_LIBRARY_ARCHITECTURE "@CMAKE_HIP_LIBRARY_ARCHITECTURE@")
+endif()
+
+set(CMAKE_HIP_TOOLKIT_INCLUDE_DIRECTORIES "@CMAKE_HIP_TOOLKIT_INCLUDE_DIRECTORIES@")
+
+set(CMAKE_HIP_IMPLICIT_INCLUDE_DIRECTORIES "@CMAKE_HIP_IMPLICIT_INCLUDE_DIRECTORIES@")
+set(CMAKE_HIP_IMPLICIT_LINK_LIBRARIES "@CMAKE_HIP_IMPLICIT_LINK_LIBRARIES@")
+set(CMAKE_HIP_IMPLICIT_LINK_DIRECTORIES "@CMAKE_HIP_IMPLICIT_LINK_DIRECTORIES@")
+set(CMAKE_HIP_IMPLICIT_LINK_FRAMEWORK_DIRECTORIES "@CMAKE_HIP_IMPLICIT_LINK_FRAMEWORK_DIRECTORIES@")
+
+set(CMAKE_HIP_COMPILER_DEVICE_LIBRARY_ROOT_DIR "@CMAKE_HIP_COMPILER_DEVICE_LIBRARY_ROOT_DIR@")
+@_SET_CMAKE_HIP_DEVICE_LIBRARY_INCLUSION@
+
+set(CMAKE_AR "@CMAKE_AR@")
+set(CMAKE_HIP_COMPILER_AR "@CMAKE_HIP_COMPILER_AR@")
+set(CMAKE_RANLIB "@CMAKE_RANLIB@")
+set(CMAKE_HIP_COMPILER_RANLIB "@CMAKE_HIP_COMPILER_RANLIB@")
+set(CMAKE_LINKER "@CMAKE_LINKER@")
+set(CMAKE_MT "@CMAKE_MT@")

Modules/CMakeHIPCompilerABI.hip

@@ -0,0 +1,16 @@
+#ifndef __HIP__
+#  error "A C or C++ compiler has been selected for HIP"
+#endif
+
+#include "CMakeCompilerABI.h"
+
+int main(int argc, char* argv[])
+{
+  int require = 0;
+  require += info_sizeof_dptr[argc];
+#if defined(ABI_ID)
+  require += info_abi[argc];
+#endif
+  (void)argv;
+  return require;
+}

Modules/CMakeHIPCompilerId.hip.in

@@ -0,0 +1,54 @@
+#ifndef __HIP__
+# error "A C or C++ compiler has been selected for HIP"
+#endif
+
+@CMAKE_HIP_COMPILER_ID_CONTENT@
+
+/* Construct the string literal in pieces to prevent the source from
+   getting matched.  Store it in a pointer rather than an array
+   because some compilers will just produce instructions to fill the
+   array rather than assigning a pointer to a static array.  */
+char const* info_compiler = "INFO" ":" "compiler[" COMPILER_ID "]";
+#ifdef SIMULATE_ID
+char const* info_simulate = "INFO" ":" "simulate[" SIMULATE_ID "]";
+#endif
+
+@CMAKE_HIP_COMPILER_ID_PLATFORM_CONTENT@
+@CMAKE_HIP_COMPILER_ID_ERROR_FOR_TEST@
+
+const char* info_language_dialect_default = "INFO" ":" "dialect_default["
+#if __cplusplus > 202002L
+  "23"
+#elif __cplusplus > 201703L
+  "20"
+#elif __cplusplus >= 201703L
+  "17"
+#elif __cplusplus >= 201402L
+  "14"
+#elif __cplusplus >= 201103L
+  "11"
+#else
+  "98"
+#endif
+"]";
+
+/*--------------------------------------------------------------------------*/
+
+int main(int argc, char* argv[])
+{
+  int require = 0;
+  require += info_compiler[argc];
+  require += info_platform[argc];
+#ifdef COMPILER_VERSION_MAJOR
+  require += info_version[argc];
+#endif
+#ifdef SIMULATE_ID
+  require += info_simulate[argc];
+#endif
+#ifdef SIMULATE_VERSION_MAJOR
+  require += info_simulate_version[argc];
+#endif
+  require += info_language_dialect_default[argc];
+  (void)argv;
+  return require;
+}

Modules/CMakeHIPInformation.cmake

@@ -0,0 +1,132 @@
+# Distributed under the OSI-approved BSD 3-Clause License.  See accompanying
+# file Copyright.txt or https://cmake.org/licensing for details.
+
+if(UNIX)
+  set(CMAKE_HIP_OUTPUT_EXTENSION .o)
+else()
+  set(CMAKE_HIP_OUTPUT_EXTENSION .obj)
+endif()
+set(CMAKE_INCLUDE_FLAG_HIP "-I")
+
+# Load compiler-specific information.
+if(CMAKE_HIP_COMPILER_ID)
+  include(Compiler/${CMAKE_HIP_COMPILER_ID}-HIP OPTIONAL)
+endif()
+
+# load the system- and compiler specific files
+if(CMAKE_HIP_COMPILER_ID)
+  # load a hardware specific file, mostly useful for embedded compilers
+  if(CMAKE_SYSTEM_PROCESSOR)
+    include(Platform/${CMAKE_EFFECTIVE_SYSTEM_NAME}-${CMAKE_HIP_COMPILER_ID}-HIP-${CMAKE_SYSTEM_PROCESSOR} OPTIONAL)
+  endif()
+  include(Platform/${CMAKE_EFFECTIVE_SYSTEM_NAME}-${CMAKE_HIP_COMPILER_ID}-HIP OPTIONAL)
+endif()
+
+
+if(NOT CMAKE_SHARED_LIBRARY_RUNTIME_HIP_FLAG)
+  set(CMAKE_SHARED_LIBRARY_RUNTIME_HIP_FLAG ${CMAKE_SHARED_LIBRARY_RUNTIME_C_FLAG})
+endif()
+
+if(NOT CMAKE_SHARED_LIBRARY_RUNTIME_HIP_FLAG_SEP)
+  set(CMAKE_SHARED_LIBRARY_RUNTIME_HIP_FLAG_SEP ${CMAKE_SHARED_LIBRARY_RUNTIME_C_FLAG_SEP})
+endif()
+
+if(NOT CMAKE_SHARED_LIBRARY_RPATH_LINK_HIP_FLAG)
+  set(CMAKE_SHARED_LIBRARY_RPATH_LINK_HIP_FLAG ${CMAKE_SHARED_LIBRARY_RPATH_LINK_C_FLAG})
+endif()
+
+if(NOT DEFINED CMAKE_EXE_EXPORTS_HIP_FLAG)
+  set(CMAKE_EXE_EXPORTS_HIP_FLAG ${CMAKE_EXE_EXPORTS_C_FLAG})
+endif()
+
+if(NOT DEFINED CMAKE_SHARED_LIBRARY_SONAME_HIP_FLAG)
+  set(CMAKE_SHARED_LIBRARY_SONAME_HIP_FLAG ${CMAKE_SHARED_LIBRARY_SONAME_C_FLAG})
+endif()
+
+if(NOT CMAKE_EXECUTABLE_RUNTIME_HIP_FLAG)
+  set(CMAKE_EXECUTABLE_RUNTIME_HIP_FLAG ${CMAKE_SHARED_LIBRARY_RUNTIME_HIP_FLAG})
+endif()
+
+if(NOT CMAKE_EXECUTABLE_RUNTIME_HIP_FLAG_SEP)
+  set(CMAKE_EXECUTABLE_RUNTIME_HIP_FLAG_SEP ${CMAKE_SHARED_LIBRARY_RUNTIME_HIP_FLAG_SEP})
+endif()
+
+if(NOT CMAKE_EXECUTABLE_RPATH_LINK_HIP_FLAG)
+  set(CMAKE_EXECUTABLE_RPATH_LINK_HIP_FLAG ${CMAKE_SHARED_LIBRARY_RPATH_LINK_HIP_FLAG})
+endif()
+
+if(NOT DEFINED CMAKE_SHARED_LIBRARY_LINK_HIP_WITH_RUNTIME_PATH)
+  set(CMAKE_SHARED_LIBRARY_LINK_HIP_WITH_RUNTIME_PATH ${CMAKE_SHARED_LIBRARY_LINK_C_WITH_RUNTIME_PATH})
+endif()
+
+
+# for most systems a module is the same as a shared library
+# so unless the variable CMAKE_MODULE_EXISTS is set just
+# copy the values from the LIBRARY variables
+if(NOT CMAKE_MODULE_EXISTS)
+  set(CMAKE_SHARED_MODULE_HIP_FLAGS ${CMAKE_SHARED_LIBRARY_HIP_FLAGS})
+  set(CMAKE_SHARED_MODULE_CREATE_HIP_FLAGS ${CMAKE_SHARED_LIBRARY_CREATE_HIP_FLAGS})
+endif()
+
+# add the flags to the cache based
+# on the initial values computed in the platform/*.cmake files
+# use _INIT variables so that this only happens the first time
+# and you can set these flags in the cmake cache
+set(CMAKE_HIP_FLAGS_INIT "$ENV{HIPFLAGS} ${CMAKE_HIP_FLAGS_INIT}")
+
+cmake_initialize_per_config_variable(CMAKE_HIP_FLAGS "Flags used by the HIP compiler")
+
+if(CMAKE_HIP_STANDARD_LIBRARIES_INIT)
+  set(CMAKE_HIP_STANDARD_LIBRARIES "${CMAKE_HIP_STANDARD_LIBRARIES_INIT}"
+    CACHE STRING "Libraries linked by default with all HIP applications.")
+  mark_as_advanced(CMAKE_HIP_STANDARD_LIBRARIES)
+endif()
+
+if(NOT CMAKE_HIP_COMPILER_LAUNCHER AND DEFINED ENV{CMAKE_HIP_COMPILER_LAUNCHER})
+  set(CMAKE_HIP_COMPILER_LAUNCHER "$ENV{CMAKE_HIP_COMPILER_LAUNCHER}"
+    CACHE STRING "Compiler launcher for HIP.")
+endif()
+
+include(CMakeCommonLanguageInclude)
+
+# now define the following rules:
+# CMAKE_HIP_CREATE_SHARED_LIBRARY
+# CMAKE_HIP_CREATE_SHARED_MODULE
+# CMAKE_HIP_COMPILE_OBJECT
+# CMAKE_HIP_LINK_EXECUTABLE
+
+# create a shared library
+if(NOT CMAKE_HIP_CREATE_SHARED_LIBRARY)
+  set(CMAKE_HIP_CREATE_SHARED_LIBRARY
+      "<CMAKE_HIP_COMPILER> <CMAKE_SHARED_LIBRARY_HIP_FLAGS> <LANGUAGE_COMPILE_FLAGS> <LINK_FLAGS> <CMAKE_SHARED_LIBRARY_CREATE_HIP_FLAGS> <SONAME_FLAG><TARGET_SONAME> -o <TARGET> <OBJECTS> <LINK_LIBRARIES>")
+endif()
+
+# create a shared module copy the shared library rule by default
+if(NOT CMAKE_HIP_CREATE_SHARED_MODULE)
+  set(CMAKE_HIP_CREATE_SHARED_MODULE ${CMAKE_HIP_CREATE_SHARED_LIBRARY})
+endif()
+
+# Create a static archive incrementally for large object file counts.
+if(NOT DEFINED CMAKE_HIP_ARCHIVE_CREATE)
+  set(CMAKE_HIP_ARCHIVE_CREATE "<CMAKE_AR> qc <TARGET> <LINK_FLAGS> <OBJECTS>")
+endif()
+if(NOT DEFINED CMAKE_HIP_ARCHIVE_APPEND)
+  set(CMAKE_HIP_ARCHIVE_APPEND "<CMAKE_AR> q <TARGET> <LINK_FLAGS> <OBJECTS>")
+endif()
+if(NOT DEFINED CMAKE_HIP_ARCHIVE_FINISH)
+  set(CMAKE_HIP_ARCHIVE_FINISH "<CMAKE_RANLIB> <TARGET>")
+endif()
+
+# compile a HIP file into an object file
+if(NOT CMAKE_HIP_COMPILE_OBJECT)
+  set(CMAKE_HIP_COMPILE_OBJECT
+    "<CMAKE_HIP_COMPILER> <DEFINES> <INCLUDES> <FLAGS> -o <OBJECT> ${_CMAKE_COMPILE_AS_HIP_FLAG} -c <SOURCE>")
+endif()
+
+# compile a cu file into an executable
+if(NOT CMAKE_HIP_LINK_EXECUTABLE)
+  set(CMAKE_HIP_LINK_EXECUTABLE
+    "<CMAKE_HIP_COMPILER> <FLAGS> <CMAKE_HIP_LINK_FLAGS> <LINK_FLAGS> <OBJECTS> -o <TARGET> <LINK_LIBRARIES>")
+endif()
+
+set(CMAKE_HIP_INFORMATION_LOADED 1)

Modules/CMakeTestHIPCompiler.cmake

@@ -0,0 +1,98 @@
+# Distributed under the OSI-approved BSD 3-Clause License.  See accompanying
+# file Copyright.txt or https://cmake.org/licensing for details.
+
+
+if(CMAKE_HIP_COMPILER_FORCED)
+  # The compiler configuration was forced by the user.
+  # Assume the user has configured all compiler information.
+  set(CMAKE_HIP_COMPILER_WORKS TRUE)
+  return()
+endif()
+
+set(__CMAKE_HIP_FLAGS "${CMAKE_HIP_FLAGS}")
+string(APPEND CMAKE_HIP_FLAGS "--cuda-host-only")
+
+include(CMakeTestCompilerCommon)
+
+# work around enforced code signing and / or missing executable target type
+set(__CMAKE_SAVED_TRY_COMPILE_TARGET_TYPE ${CMAKE_TRY_COMPILE_TARGET_TYPE})
+if(_CMAKE_FEATURE_DETECTION_TARGET_TYPE)
+  set(CMAKE_TRY_COMPILE_TARGET_TYPE ${_CMAKE_FEATURE_DETECTION_TARGET_TYPE})
+endif()
+
+# Remove any cached result from an older CMake version.
+# We now store this in CMakeHIPCompiler.cmake.
+unset(CMAKE_HIP_COMPILER_WORKS CACHE)
+
+# Try to identify the ABI and configure it into CMakeHIPCompiler.cmake
+include(${CMAKE_ROOT}/Modules/CMakeDetermineCompilerABI.cmake)
+CMAKE_DETERMINE_COMPILER_ABI(HIP ${CMAKE_ROOT}/Modules/CMakeHIPCompilerABI.hip)
+if(CMAKE_HIP_ABI_COMPILED)
+  # The compiler worked so skip dedicated test below.
+  set(CMAKE_HIP_COMPILER_WORKS TRUE)
+  message(STATUS "Check for working HIP compiler: ${CMAKE_HIP_COMPILER} - skipped")
+endif()
+
+# This file is used by EnableLanguage in cmGlobalGenerator to
+# determine that the selected C++ compiler can actually compile
+# and link the most basic of programs.   If not, a fatal error
+# is set and cmake stops processing commands and will not generate
+# any makefiles or projects.
+if(NOT CMAKE_HIP_COMPILER_WORKS)
+  PrintTestCompilerStatus("HIP")
+  __TestCompiler_setTryCompileTargetType()
+  file(WRITE ${CMAKE_BINARY_DIR}${CMAKE_FILES_DIRECTORY}/CMakeTmp/testHIPCompiler.hip
+    "#ifndef __HIP__\n"
+    "# error \"The CMAKE_HIP_COMPILER is set to a C/CXX compiler\"\n"
+    "#endif\n"
+    "int main(){return 0;}\n")
+  try_compile(CMAKE_HIP_COMPILER_WORKS ${CMAKE_BINARY_DIR}
+    ${CMAKE_BINARY_DIR}${CMAKE_FILES_DIRECTORY}/CMakeTmp/testHIPCompiler.hip
+    OUTPUT_VARIABLE __CMAKE_HIP_COMPILER_OUTPUT)
+  # Move result from cache to normal variable.
+  set(CMAKE_HIP_COMPILER_WORKS ${CMAKE_HIP_COMPILER_WORKS})
+  unset(CMAKE_HIP_COMPILER_WORKS CACHE)
+  __TestCompiler_restoreTryCompileTargetType()
+  if(NOT CMAKE_HIP_COMPILER_WORKS)
+    PrintTestCompilerResult(CHECK_FAIL "broken")
+    file(APPEND ${CMAKE_BINARY_DIR}${CMAKE_FILES_DIRECTORY}/CMakeError.log
+      "Determining if the HIP compiler works failed with "
+      "the following output:\n${__CMAKE_HIP_COMPILER_OUTPUT}\n\n")
+    string(REPLACE "\n" "\n  " _output "${__CMAKE_HIP_COMPILER_OUTPUT}")
+    message(FATAL_ERROR "The HIP compiler\n  \"${CMAKE_HIP_COMPILER}\"\n"
+      "is not able to compile a simple test program.\nIt fails "
+      "with the following output:\n  ${_output}\n\n"
+      "CMake will not be able to correctly generate this project.")
+  endif()
+  PrintTestCompilerResult(CHECK_PASS "works")
+  file(APPEND ${CMAKE_BINARY_DIR}${CMAKE_FILES_DIRECTORY}/CMakeOutput.log
+    "Determining if the HIP compiler works passed with "
+    "the following output:\n${__CMAKE_HIP_COMPILER_OUTPUT}\n\n")
+endif()
+
+set(CMAKE_HIP_FLAGS "${__CMAKE_HIP_FLAGS}")
+unset(__CMAKE_HIP_FLAGS)
+
+
+# Try to identify the compiler features
+include(${CMAKE_ROOT}/Modules/CMakeDetermineCompileFeatures.cmake)
+CMAKE_DETERMINE_COMPILE_FEATURES(HIP)
+
+# Re-configure to save learned information.
+configure_file(
+  ${CMAKE_ROOT}/Modules/CMakeHIPCompiler.cmake.in
+  ${CMAKE_PLATFORM_INFO_DIR}/CMakeHIPCompiler.cmake
+  @ONLY
+  )
+include(${CMAKE_PLATFORM_INFO_DIR}/CMakeHIPCompiler.cmake)
+
+if(CMAKE_HIP_SIZEOF_DATA_PTR)
+  foreach(f ${CMAKE_HIP_ABI_FILES})
+    include(${f})
+  endforeach()
+  unset(CMAKE_HIP_ABI_FILES)
+endif()
+
+set(CMAKE_TRY_COMPILE_TARGET_TYPE ${__CMAKE_SAVED_TRY_COMPILE_TARGET_TYPE})
+unset(__CMAKE_SAVED_TRY_COMPILE_TARGET_TYPE)
+unset(__CMAKE_HIP_COMPILER_OUTPUT)

Modules/Compiler/CMakeCommonCompilerMacros.cmake

@@ -170,3 +170,19 @@ macro(cmake_record_cuda_compile_features)
     unset(CMAKE_CUDA03_STANDARD__HAS_FULL_SUPPORT)
   endif()
 endmacro()
+
+macro(cmake_record_hip_compile_features)
+  set(_result 0)
+  if(_result EQUAL 0 AND DEFINED CMAKE_HIP23_STANDARD_COMPILE_OPTION)
+    _has_compiler_features_hip(23)
+  endif()
+  if(_result EQUAL 0 AND DEFINED CMAKE_HIP20_STANDARD_COMPILE_OPTION)
+    _has_compiler_features_hip(20)
+  endif()
+  if(_result EQUAL 0 AND DEFINED CMAKE_HIP17_STANDARD_COMPILE_OPTION)
+    _has_compiler_features_hip(17)
+  endif()
+  _has_compiler_features_hip(14)
+  _has_compiler_features_hip(11)
+  _has_compiler_features_hip(98)
+endmacro()

Modules/Compiler/Clang-HIP.cmake

@@ -0,0 +1,16 @@
+include(Compiler/Clang)
+__compiler_clang(HIP)
+__compiler_clang_cxx_standards(HIP)
+
+set(_CMAKE_COMPILE_AS_HIP_FLAG "-x hip")
+set(_CMAKE_HIP_RDC_FLAG "-fgpu-rdc")
+
+if(NOT "x${CMAKE_HIP_SIMULATE_ID}" STREQUAL "xMSVC")
+  set(CMAKE_HIP_COMPILE_OPTIONS_VISIBILITY_INLINES_HIDDEN "-fvisibility-inlines-hidden")
+
+  string(APPEND CMAKE_HIP_FLAGS_DEBUG_INIT " -O")
+endif()
+
+set(CMAKE_HIP_RUNTIME_LIBRARY_DEFAULT "SHARED")
+set(CMAKE_HIP_RUNTIME_LIBRARY_LINK_OPTIONS_STATIC  "")
+set(CMAKE_HIP_RUNTIME_LIBRARY_LINK_OPTIONS_SHARED  "")

Modules/Compiler/ROCMClang-HIP.cmake

@@ -0,0 +1,45 @@
+include(Compiler/ROCMClang)
+__compiler_rocmclang(HIP)
+
+set(_CMAKE_COMPILE_AS_HIP_FLAG "-x hip")
+set(_CMAKE_HIP_RDC_FLAG "-fgpu-rdc")
+
+if(NOT "x${CMAKE_${lang}_SIMULATE_ID}" STREQUAL "xMSVC")
+  set(CMAKE_HIP_COMPILE_OPTIONS_VISIBILITY_INLINES_HIDDEN "-fvisibility-inlines-hidden")
+  string(APPEND CMAKE_HIP_FLAGS_DEBUG_INIT " -O")
+endif()
+
+if(CMAKE_HIP_SIMULATE_ID STREQUAL "GNU")
+  set(CMAKE_HIP_LINKER_WRAPPER_FLAG "-Wl,")
+  set(CMAKE_HIP_LINKER_WRAPPER_FLAG_SEP ",")
+elseif(CMAKE_HIP_SIMULATE_ID STREQUAL "Clang")
+  set(CMAKE_HIP_LINKER_WRAPPER_FLAG "-Xlinker" " ")
+  set(CMAKE_HIP_LINKER_WRAPPER_FLAG_SEP)
+endif()
+
+if(NOT CMAKE_HIP_COMPILER_VERSION VERSION_LESS 1.0)
+  set(CMAKE_HIP98_STANDARD_COMPILE_OPTION "-std=c++98")
+  set(CMAKE_HIP98_EXTENSION_COMPILE_OPTION "-std=gnu++98")
+  set(CMAKE_HIP98_STANDARD__HAS_FULL_SUPPORT ON)
+
+  set(CMAKE_HIP11_STANDARD_COMPILE_OPTION "-std=c++11")
+  set(CMAKE_HIP11_EXTENSION_COMPILE_OPTION "-std=gnu++11")
+  set(CMAKE_HIP11_STANDARD__HAS_FULL_SUPPORT ON)
+
+  set(CMAKE_HIP14_STANDARD_COMPILE_OPTION "-std=c++14")
+  set(CMAKE_HIP14_EXTENSION_COMPILE_OPTION "-std=gnu++14")
+  set(CMAKE_HIP14_STANDARD__HAS_FULL_SUPPORT ON)
+
+  set(CMAKE_HIP17_STANDARD_COMPILE_OPTION "-std=c++17")
+  set(CMAKE_HIP17_EXTENSION_COMPILE_OPTION "-std=gnu++17")
+  set(CMAKE_HIP17_STANDARD__HAS_FULL_SUPPORT ON)
+
+  set(CMAKE_HIP20_STANDARD_COMPILE_OPTION "-std=c++20")
+  set(CMAKE_HIP20_EXTENSION_COMPILE_OPTION "-std=gnu++20")
+endif()
+
+set(CMAKE_HIP_RUNTIME_LIBRARY_DEFAULT "SHARED")
+set(CMAKE_HIP_RUNTIME_LIBRARY_LINK_OPTIONS_STATIC  "")
+set(CMAKE_HIP_RUNTIME_LIBRARY_LINK_OPTIONS_SHARED  "")
+
+__compiler_check_default_language_standard(HIP 3.5 11)

Modules/Internal/FeatureTesting.cmake

@@ -99,6 +99,16 @@ macro(_record_compiler_features_cuda std)
   unset(lang_level_has_features)
 endmacro()
 
+macro(_record_compiler_features_hip std)
+  list(APPEND CMAKE_HIP${std}_COMPILE_FEATURES hip_std_${std})
+
+  get_property(lang_level_has_features GLOBAL PROPERTY CMAKE_HIP${std}_KNOWN_FEATURES)
+  if(lang_level_has_features)
+    _record_compiler_features(HIP "${CMAKE_HIP${std}_STANDARD_COMPILE_OPTION}" CMAKE_HIP${std}_COMPILE_FEATURES)
+  endif()
+  unset(lang_level_has_features)
+endmacro()
+
 macro(_has_compiler_features lang level compile_flags feature_list)
   # presume all known features are supported
   get_property(known_features GLOBAL PROPERTY CMAKE_${lang}${level}_KNOWN_FEATURES)
@@ -117,3 +127,7 @@ macro(_has_compiler_features_cuda std)
   list(APPEND CMAKE_CUDA${std}_COMPILE_FEATURES cuda_std_${std})
   _has_compiler_features(CUDA ${std} "${CMAKE_CUDA${std}_STANDARD_COMPILE_OPTION}" CMAKE_CUDA${std}_COMPILE_FEATURES)
 endmacro()
+macro(_has_compiler_features_hip std)
+  list(APPEND CMAKE_HIP${std}_COMPILE_FEATURES hip_std_${std})
+  _has_compiler_features(HIP ${std} "${CMAKE_HIP${std}_STANDARD_COMPILE_OPTION}" CMAKE_HIP${std}_COMPILE_FEATURES)
+endmacro()

Source/cmComputeLinkInformation.cxx

@@ -585,10 +585,10 @@ void cmComputeLinkInformation::AddImplicitLinkInfo()
     this->Target->GetLinkClosure(this->Config);
   for (std::string const& li : lc->Languages) {
 
-    if (li == "CUDA") {
+    if (li == "CUDA" || li == "HIP") {
       // These need to go before the other implicit link information
       // as they could require symbols from those other library
-      // Currently restricted to CUDA as it is the only language
+      // Currently restricted as CUDA and HIP are the only languages
       // we have documented runtime behavior controls for
       this->AddRuntimeLinkLibrary(li);
     }

Source/cmCoreTryCompile.cxx

@@ -189,6 +189,8 @@ SETUP_LANGUAGE(cuda_properties, CUDA);
 // NOLINTNEXTLINE(bugprone-suspicious-missing-comma)
 SETUP_LANGUAGE(fortran_properties, Fortran);
 // NOLINTNEXTLINE(bugprone-suspicious-missing-comma)
+SETUP_LANGUAGE(hip_properties, HIP);
+// NOLINTNEXTLINE(bugprone-suspicious-missing-comma)
 SETUP_LANGUAGE(objc_properties, OBJC);
 // NOLINTNEXTLINE(bugprone-suspicious-missing-comma)
 SETUP_LANGUAGE(objcxx_properties, OBJCXX);
@@ -201,6 +203,8 @@ SETUP_LANGUAGE(swift_properties, Swift);
 std::string const kCMAKE_CUDA_ARCHITECTURES = "CMAKE_CUDA_ARCHITECTURES";
 std::string const kCMAKE_CUDA_RUNTIME_LIBRARY = "CMAKE_CUDA_RUNTIME_LIBRARY";
 std::string const kCMAKE_ENABLE_EXPORTS = "CMAKE_ENABLE_EXPORTS";
+std::string const kCMAKE_HIP_ARCHITECTURES = "CMAKE_HIP_ARCHITECTURES";
+std::string const kCMAKE_HIP_RUNTIME_LIBRARY = "CMAKE_HIP_RUNTIME_LIBRARY";
 std::string const kCMAKE_ISPC_INSTRUCTION_SETS = "CMAKE_ISPC_INSTRUCTION_SETS";
 std::string const kCMAKE_ISPC_HEADER_SUFFIX = "CMAKE_ISPC_HEADER_SUFFIX";
 std::string const kCMAKE_LINK_SEARCH_END_STATIC =
@@ -274,6 +278,7 @@ int cmCoreTryCompile::TryCompileCode(std::vector<std::string> const& argv,
   LanguageStandardState cState("C");
   LanguageStandardState cudaState("CUDA");
   LanguageStandardState cxxState("CXX");
+  LanguageStandardState hipState("HIP");
   LanguageStandardState objcState("OBJC");
   LanguageStandardState objcxxState("OBJCXX");
   std::vector<std::string> targets;
@@ -323,6 +328,7 @@ int cmCoreTryCompile::TryCompileCode(std::vector<std::string> const& argv,
     } else if (cState.UpdateIfMatches(argv, i) ||
                cxxState.UpdateIfMatches(argv, i) ||
                cudaState.UpdateIfMatches(argv, i) ||
+               hipState.UpdateIfMatches(argv, i) ||
                objcState.UpdateIfMatches(argv, i) ||
                objcxxState.UpdateIfMatches(argv, i)) {
       continue;
@@ -428,6 +434,9 @@ int cmCoreTryCompile::TryCompileCode(std::vector<std::string> const& argv,
     if (!cudaState.Validate(this->Makefile)) {
       return -1;
     }
+    if (!hipState.Validate(this->Makefile)) {
+      return -1;
+    }
     if (!cxxState.Validate(this->Makefile)) {
       return -1;
     }
@@ -715,6 +724,8 @@ int cmCoreTryCompile::TryCompileCode(std::vector<std::string> const& argv,
       vars.insert(
         &fortran_properties[lang_property_start],
         &fortran_properties[lang_property_start + lang_property_size]);
+      vars.insert(&hip_properties[lang_property_start],
+                  &hip_properties[lang_property_start + lang_property_size]);
       vars.insert(&objc_properties[lang_property_start],
                   &objc_properties[lang_property_start + lang_property_size]);
       vars.insert(
@@ -727,6 +738,8 @@ int cmCoreTryCompile::TryCompileCode(std::vector<std::string> const& argv,
       vars.insert(kCMAKE_CUDA_ARCHITECTURES);
       vars.insert(kCMAKE_CUDA_RUNTIME_LIBRARY);
       vars.insert(kCMAKE_ENABLE_EXPORTS);
+      vars.insert(kCMAKE_HIP_ARCHITECTURES);
+      vars.insert(kCMAKE_HIP_RUNTIME_LIBRARY);
       vars.insert(kCMAKE_ISPC_INSTRUCTION_SETS);
       vars.insert(kCMAKE_ISPC_HEADER_SUFFIX);
       vars.insert(kCMAKE_LINK_SEARCH_END_STATIC);
@@ -761,6 +774,8 @@ int cmCoreTryCompile::TryCompileCode(std::vector<std::string> const& argv,
         vars.insert(
           &fortran_properties[pie_property_start],
           &fortran_properties[pie_property_start + pie_property_size]);
+        vars.insert(&hip_properties[pie_property_start],
+                    &hip_properties[pie_property_start + pie_property_size]);
         vars.insert(&objc_properties[pie_property_start],
                     &objc_properties[pie_property_start + pie_property_size]);
         vars.insert(
@@ -835,6 +850,7 @@ int cmCoreTryCompile::TryCompileCode(std::vector<std::string> const& argv,
     cState.Enabled(testLangs.find("C") != testLangs.end());
     cxxState.Enabled(testLangs.find("CXX") != testLangs.end());
     cudaState.Enabled(testLangs.find("CUDA") != testLangs.end());
+    hipState.Enabled(testLangs.find("HIP") != testLangs.end());
     objcState.Enabled(testLangs.find("OBJC") != testLangs.end());
     objcxxState.Enabled(testLangs.find("OBJCXX") != testLangs.end());
 
@@ -842,7 +858,7 @@ int cmCoreTryCompile::TryCompileCode(std::vector<std::string> const& argv,
     bool honorStandard = true;
 
     if (cState.DidNone() && cxxState.DidNone() && objcState.DidNone() &&
-        objcxxState.DidNone() && cudaState.DidNone()) {
+        objcxxState.DidNone() && cudaState.DidNone() && hipState.DidNone()) {
       switch (this->Makefile->GetPolicyStatus(cmPolicies::CMP0067)) {
         case cmPolicies::WARN:
           warnCMP0067 = this->Makefile->PolicyOptionalWarningEnabled(
@@ -872,6 +888,8 @@ int cmCoreTryCompile::TryCompileCode(std::vector<std::string> const& argv,
                                      warnCMP0067, warnCMP0067Variables);
     cudaState.LoadUnsetPropertyValues(this->Makefile, honorStandard,
                                       warnCMP0067, warnCMP0067Variables);
+    hipState.LoadUnsetPropertyValues(this->Makefile, honorStandard,
+                                     warnCMP0067, warnCMP0067Variables);
     objcState.LoadUnsetPropertyValues(this->Makefile, honorStandard,
                                       warnCMP0067, warnCMP0067Variables);
     objcxxState.LoadUnsetPropertyValues(this->Makefile, honorStandard,
@@ -894,6 +912,7 @@ int cmCoreTryCompile::TryCompileCode(std::vector<std::string> const& argv,
     cState.WriteProperties(fout, targetName);
     cxxState.WriteProperties(fout, targetName);
     cudaState.WriteProperties(fout, targetName);
+    hipState.WriteProperties(fout, targetName);
     objcState.WriteProperties(fout, targetName);
     objcxxState.WriteProperties(fout, targetName);
 

Source/cmExtraCodeBlocksGenerator.cxx

@@ -366,10 +366,11 @@ void cmExtraCodeBlocksGenerator::CreateNewProjectFile(
               continue;
             }
 
-            // check whether it is a C/C++/CUDA implementation file
+            // check whether it is a C/C++/CUDA/HIP implementation file
             bool isCFile = false;
             std::string lang = s->GetOrDetermineLanguage();
-            if (lang == "C" || lang == "CXX" || lang == "CUDA") {
+            if (lang == "C" || lang == "CXX" || lang == "CUDA" ||
+                lang == "HIP") {
               std::string const& srcext = s->GetExtension();
               isCFile = cm->IsACLikeSourceExtension(srcext);
             }

Source/cmGeneratorExpressionNode.cxx

@@ -718,7 +718,7 @@ struct CompilerIdNode : public cmGeneratorExpressionNode
 static const CompilerIdNode cCompilerIdNode("C"), cxxCompilerIdNode("CXX"),
   cudaCompilerIdNode("CUDA"), objcCompilerIdNode("OBJC"),
   objcxxCompilerIdNode("OBJCXX"), fortranCompilerIdNode("Fortran"),
-  ispcCompilerIdNode("ISPC");
+  hipCompilerIdNode("HIP"), ispcCompilerIdNode("ISPC");
 
 struct CompilerVersionNode : public cmGeneratorExpressionNode
 {
@@ -783,7 +783,8 @@ struct CompilerVersionNode : public cmGeneratorExpressionNode
 static const CompilerVersionNode cCompilerVersionNode("C"),
   cxxCompilerVersionNode("CXX"), cudaCompilerVersionNode("CUDA"),
   objcCompilerVersionNode("OBJC"), objcxxCompilerVersionNode("OBJCXX"),
-  fortranCompilerVersionNode("Fortran"), ispcCompilerVersionNode("ISPC");
+  fortranCompilerVersionNode("Fortran"), ispcCompilerVersionNode("ISPC"),
+  hipCompilerVersionNode("HIP");
 
 struct PlatformIdNode : public cmGeneratorExpressionNode
 {
@@ -2597,6 +2598,7 @@ const cmGeneratorExpressionNode* cmGeneratorExpressionNode::GetNode(
     { "OBJCXX_COMPILER_ID", &objcxxCompilerIdNode },
     { "CUDA_COMPILER_ID", &cudaCompilerIdNode },
     { "Fortran_COMPILER_ID", &fortranCompilerIdNode },
+    { "HIP_COMPILER_ID", &hipCompilerIdNode },
     { "VERSION_GREATER", &versionGreaterNode },
     { "VERSION_GREATER_EQUAL", &versionGreaterEqNode },
     { "VERSION_LESS", &versionLessNode },
@@ -2608,6 +2610,7 @@ const cmGeneratorExpressionNode* cmGeneratorExpressionNode::GetNode(
     { "OBJC_COMPILER_VERSION", &objcCompilerVersionNode },
     { "OBJCXX_COMPILER_VERSION", &objcxxCompilerVersionNode },
     { "Fortran_COMPILER_VERSION", &fortranCompilerVersionNode },
+    { "HIP_COMPILER_VERSION", &hipCompilerVersionNode },
     { "PLATFORM_ID", &platformIdNode },
     { "COMPILE_FEATURES", &compileFeaturesNode },
     { "CONFIGURATION", &configurationNode },

Source/cmGeneratorTarget.cxx

@@ -980,7 +980,7 @@ cmProp cmGeneratorTarget::GetPropertyWithPairedLanguageSupport(
     // Check if we should use the value set by another language.
     if (lang == "OBJC") {
       propertyValue = this->GetPropertyWithPairedLanguageSupport("C", suffix);
-    } else if (lang == "OBJCXX" || lang == "CUDA") {
+    } else if (lang == "OBJCXX" || lang == "CUDA" || lang == "HIP") {
       propertyValue =
         this->GetPropertyWithPairedLanguageSupport("CXX", suffix);
     }
@@ -1121,7 +1121,7 @@ void cmGeneratorTarget::AppendLanguageSideEffects(
   std::map<std::string, std::set<cmGeneratorTarget const*>>& sideEffects) const
 {
   static const std::set<cm::string_view> LANGS_WITH_NO_SIDE_EFFECTS = {
-    "C"_s, "CXX"_s, "OBJC"_s, "OBJCXX"_s, "ASM"_s, "CUDA"_s,
+    "C"_s, "CXX"_s, "OBJC"_s, "OBJCXX"_s, "ASM"_s, "CUDA"_s, "HIP"_s
   };
 
   for (auto const& lang : this->GetAllConfigCompileLanguages()) {
@@ -3351,6 +3351,23 @@ void cmGeneratorTarget::AddISPCTargetFlags(std::string& flags) const
   }
 }
 
+void cmGeneratorTarget::AddHIPArchitectureFlags(std::string& flags) const
+{
+  const std::string& property = this->GetSafeProperty("HIP_ARCHITECTURES");
+
+  // If HIP_ARCHITECTURES is false we don't add any architectures.
+  if (cmIsOff(property)) {
+    return;
+  }
+
+  std::vector<std::string> options;
+  cmExpandList(property, options);
+
+  for (std::string& option : options) {
+    flags += " --offload-arch=" + option;
+  }
+}
+
 void cmGeneratorTarget::AddCUDAToolkitFlags(std::string& flags) const
 {
   std::string const& compiler =
@@ -4742,7 +4759,8 @@ bool cmGeneratorTarget::ComputeCompileFeatures(
       }
 
       // Custom updates for the CUDA standard.
-      if (generatorTargetLanguageStandard && language.first == "CUDA") {
+      if (generatorTargetLanguageStandard != nullptr &&
+          (language.first == "CUDA")) {
         if (generatorTargetLanguageStandard->Value == "98") {
           this->LanguageStandardMap[key].Value = "03";
         }

Source/cmGeneratorTarget.h

@@ -459,6 +459,8 @@ public:
   void AddCUDAArchitectureFlags(std::string& flags) const;
   void AddCUDAToolkitFlags(std::string& flags) const;
 
+  void AddHIPArchitectureFlags(std::string& flags) const;
+
   void AddISPCTargetFlags(std::string& flags) const;
 
   std::string GetFeatureSpecificLinkRuleVariable(

Source/cmLocalGenerator.cxx

@@ -780,9 +780,9 @@ bool cmLocalGenerator::ComputeTargetCompileFeatures()
     this->Makefile->GetGeneratorConfigs(cmMakefile::IncludeEmptyConfig);
 
   using LanguagePair = std::pair<std::string, std::string>;
-  std::vector<LanguagePair> pairedLanguages{ { "OBJC", "C" },
-                                             { "OBJCXX", "CXX" },
-                                             { "CUDA", "CXX" } };
+  std::vector<LanguagePair> pairedLanguages{
+    { "OBJC", "C" }, { "OBJCXX", "CXX" }, { "CUDA", "CXX" }, { "HIP", "CXX" }
+  };
   std::set<LanguagePair> inferredEnabledLanguages;
   for (auto const& lang : pairedLanguages) {
     if (this->Makefile->GetState()->GetLanguageEnabled(lang.first)) {
@@ -1404,8 +1404,8 @@ void cmLocalGenerator::GetDeviceLinkFlags(
     linkLineComputer->GetLinkerLanguage(target, config);
 
   if (pcli) {
-    // Compute the required cuda device link libraries when
-    // resolving cuda device symbols
+    // Compute the required device link libraries when
+    // resolving gpu lang device symbols
     this->OutputLinkLibraries(pcli, linkLineComputer, linkLibs, frameworkPath,
                               linkPath);
   }
@@ -1968,6 +1968,8 @@ void cmLocalGenerator::AddLanguageFlags(std::string& flags,
       compilerSimulateId =
         this->Makefile->GetSafeDefinition("CMAKE_CXX_SIMULATE_ID");
     }
+  } else if (lang == "HIP") {
+    target->AddHIPArchitectureFlags(flags);
   }
 
   // Add VFS Overlay for Clang compilers

Source/cmLocalUnixMakefileGenerator3.cxx

@@ -289,9 +289,9 @@ void cmLocalUnixMakefileGenerator3::WriteLocalMakefile()
     for (LocalObjectEntry const& entry : localObjectFile.second) {
       if (entry.Language == "C" || entry.Language == "CXX" ||
           entry.Language == "CUDA" || entry.Language == "Fortran" ||
-          entry.Language == "ISPC") {
-        // Right now, C, C++, Fortran and CUDA have both a preprocessor and the
-        // ability to generate assembly code
+          entry.Language == "HIP" || entry.Language == "ISPC") {
+        // Right now, C, C++, CUDA, Fortran, HIP and ISPC have both a
+        // preprocessor and the ability to generate assembly code
         lang_has_preprocessor = true;
         lang_has_assembly = true;
         break;
@@ -1550,7 +1550,7 @@ bool cmLocalUnixMakefileGenerator3::ScanDependencies(
     std::unique_ptr<cmDepends> scanner;
     if (lang == "C" || lang == "CXX" || lang == "RC" || lang == "ASM" ||
         lang == "OBJC" || lang == "OBJCXX" || lang == "CUDA" ||
-        lang == "ISPC") {
+        lang == "HIP" || lang == "ISPC") {
       // TODO: Handle RC (resource files) dependencies correctly.
       scanner = cm::make_unique<cmDependsC>(this, targetDir, lang, &validDeps);
     }

Source/cmMakefileTargetGenerator.cxx

@@ -945,7 +945,8 @@ void cmMakefileTargetGenerator::WriteObjectRuleFiles(
     std::string compilerLauncher;
     if (!compileCommands.empty() &&
         (lang == "C" || lang == "CXX" || lang == "Fortran" || lang == "CUDA" ||
-         lang == "ISPC" || lang == "OBJC" || lang == "OBJCXX")) {
+         lang == "HIP" || lang == "ISPC" || lang == "OBJC" ||
+         lang == "OBJCXX")) {
       std::string const clauncher_prop = lang + "_COMPILER_LAUNCHER";
       cmProp clauncher = this->GeneratorTarget->GetProperty(clauncher_prop);
       if (cmNonempty(clauncher)) {

Source/cmNinjaTargetGenerator.cxx

@@ -833,7 +833,8 @@ void cmNinjaTargetGenerator::WriteCompileRule(const std::string& lang,
   std::string compilerLauncher;
   if (!compileCmds.empty() &&
       (lang == "C" || lang == "CXX" || lang == "Fortran" || lang == "CUDA" ||
-       lang == "ISPC" || lang == "OBJC" || lang == "OBJCXX")) {
+       lang == "HIP" || lang == "ISPC" || lang == "OBJC" ||
+       lang == "OBJCXX")) {
     std::string const clauncher_prop = cmStrCat(lang, "_COMPILER_LAUNCHER");
     cmProp clauncher = this->GeneratorTarget->GetProperty(clauncher_prop);
     if (cmNonempty(clauncher)) {

Source/cmStandardLevelResolver.cxx

@@ -36,6 +36,9 @@ const char* const CXX_FEATURES[] = { nullptr FOR_EACH_CXX_FEATURE(
 
 const char* const CUDA_FEATURES[] = { nullptr FOR_EACH_CUDA_FEATURE(
   FEATURE_STRING) };
+
+const char* const HIP_FEATURES[] = { nullptr FOR_EACH_HIP_FEATURE(
+  FEATURE_STRING) };
 #undef FEATURE_STRING
 
 struct StandardNeeded
@@ -306,8 +309,7 @@ struct StanardLevelComputer
 };
 
 std::unordered_map<std::string, StanardLevelComputer> StandardComputerMapping =
-  {
-    { "C",
+  { { "C",
       StanardLevelComputer{
         "C", std::vector<int>{ 90, 99, 11, 17, 23 },
         std::vector<std::string>{ "90", "99", "11", "17", "23" } } },
@@ -326,7 +328,10 @@ std::unordered_map<std::string, StanardLevelComputer> StandardComputerMapping =
       StanardLevelComputer{
         "OBJCXX", std::vector<int>{ 98, 11, 14, 17, 20, 23 },
         std::vector<std::string>{ "98", "11", "14", "17", "20", "23" } } },
-  };
+    { "HIP",
+      StanardLevelComputer{
+        "HIP", std::vector<int>{ 98, 11, 14, 17, 20, 23 },
+        std::vector<std::string>{ "98", "11", "14", "17", "20", "23" } } } };
 }
 
 std::string cmStandardLevelResolver::GetCompileOptionDef(
@@ -434,6 +439,13 @@ bool cmStandardLevelResolver::CompileFeatureKnown(
     lang = "CUDA";
     return true;
   }
+  bool isHIPFeature =
+    std::find_if(cm::cbegin(HIP_FEATURES) + 1, cm::cend(HIP_FEATURES),
+                 cmStrCmp(feature)) != cm::cend(HIP_FEATURES);
+  if (isHIPFeature) {
+    lang = "HIP";
+    return true;
+  }
   std::ostringstream e;
   if (error) {
     e << "specified";

Source/cmTarget.cxx

@@ -288,6 +288,7 @@ cmTarget::cmTarget(std::string const& name, cmStateEnums::TargetType type,
     SETUP_COMMON_LANGUAGE_PROPERTIES(CXX);
     SETUP_COMMON_LANGUAGE_PROPERTIES(OBJCXX);
     SETUP_COMMON_LANGUAGE_PROPERTIES(CUDA);
+    SETUP_COMMON_LANGUAGE_PROPERTIES(HIP);
 
     initProp("ANDROID_API");
     initProp("ANDROID_API_MIN");
@@ -365,6 +366,8 @@ cmTarget::cmTarget(std::string const& name, cmStateEnums::TargetType type,
     initProp("CUDA_RESOLVE_DEVICE_SYMBOLS");
     initProp("CUDA_RUNTIME_LIBRARY");
     initProp("CUDA_ARCHITECTURES");
+    initProp("HIP_RUNTIME_LIBRARY");
+    initProp("HIP_ARCHITECTURES");
     initProp("VISIBILITY_INLINES_HIDDEN");
     initProp("JOB_POOL_COMPILE");
     initProp("JOB_POOL_LINK");
@@ -1168,6 +1171,7 @@ void cmTarget::SetProperty(const std::string& prop, const char* value)
   MAKE_STATIC_PROP(C_STANDARD);
   MAKE_STATIC_PROP(CXX_STANDARD);
   MAKE_STATIC_PROP(CUDA_STANDARD);
+  MAKE_STATIC_PROP(HIP_STANDARD);
   MAKE_STATIC_PROP(OBJC_STANDARD);
   MAKE_STATIC_PROP(OBJCXX_STANDARD);
   MAKE_STATIC_PROP(COMPILE_DEFINITIONS);
@@ -1354,8 +1358,8 @@ void cmTarget::SetProperty(const std::string& prop, const char* value)
     this->SetProperty("COMPILE_PDB_NAME", cmToCStr(tmp));
     this->AddUtility(reusedFrom, false, this->impl->Makefile);
   } else if (prop == propC_STANDARD || prop == propCXX_STANDARD ||
-             prop == propCUDA_STANDARD || prop == propOBJC_STANDARD ||
-             prop == propOBJCXX_STANDARD) {
+             prop == propCUDA_STANDARD || prop == propHIP_STANDARD ||
+             prop == propOBJC_STANDARD || prop == propOBJCXX_STANDARD) {
     if (value) {
       this->impl->LanguageStandardProperties[prop] =
         BTs<std::string>(value, this->impl->Makefile->GetBacktrace());
@@ -1461,8 +1465,8 @@ void cmTarget::AppendProperty(const std::string& prop,
     this->impl->Makefile->IssueMessage(
       MessageType::FATAL_ERROR, prop + " property may not be APPENDed.");
   } else if (prop == "C_STANDARD" || prop == "CXX_STANDARD" ||
-             prop == "CUDA_STANDARD" || prop == "OBJC_STANDARD" ||
-             prop == "OBJCXX_STANDARD") {
+             prop == "CUDA_STANDARD" || prop == "HIP_STANDARD" ||
+             prop == "OBJC_STANDARD" || prop == "OBJCXX_STANDARD") {
     this->impl->Makefile->IssueMessage(
       MessageType::FATAL_ERROR, prop + " property may not be appended.");
   } else {

Source/cmake.cxx

@@ -218,6 +218,7 @@ cmake::cmake(Role role, cmState::Mode mode)
     setupExts(this->CudaFileExtensions, { "cu" });
     setupExts(this->FortranFileExtensions,
               { "f", "F", "for", "f77", "f90", "f95", "f03" });
+    setupExts(this->HipFileExtensions, { "hip" });
     setupExts(this->ISPCFileExtensions, { "ispc" });
   }
 }
@@ -2462,6 +2463,8 @@ std::vector<std::string> cmake::GetAllExtensions() const
   // cuda extensions are also in SourceFileExtensions so we ignore it here
   allExt.insert(allExt.end(), this->FortranFileExtensions.ordered.begin(),
                 this->FortranFileExtensions.ordered.end());
+  allExt.insert(allExt.end(), this->HipFileExtensions.ordered.begin(),
+                this->HipFileExtensions.ordered.end());
   allExt.insert(allExt.end(), this->ISPCFileExtensions.ordered.begin(),
                 this->ISPCFileExtensions.ordered.end());
   return allExt;

Source/cmake.h

@@ -297,7 +297,7 @@ public:
     return this->CLikeSourceFileExtensions.Test(ext) ||
       this->CudaFileExtensions.Test(ext) ||
       this->FortranFileExtensions.Test(ext) ||
-      this->ISPCFileExtensions.Test(ext);
+      this->HipFileExtensions.Test(ext) || this->ISPCFileExtensions.Test(ext);
   }
 
   bool IsACLikeSourceExtension(cm::string_view ext) const
@@ -662,6 +662,7 @@ private:
   FileExtensions CudaFileExtensions;
   FileExtensions ISPCFileExtensions;
   FileExtensions FortranFileExtensions;
+  FileExtensions HipFileExtensions;
   bool ClearBuildSystem = false;
   bool DebugTryCompile = false;
   bool RegenerateDuringBuild = false;
@@ -838,3 +839,11 @@ private:
   F(cuda_std_17)                                                              \
   F(cuda_std_20)                                                              \
   F(cuda_std_23)
+
+#define FOR_EACH_HIP_FEATURE(F)                                               \
+  F(hip_std_98)                                                               \
+  F(hip_std_11)                                                               \
+  F(hip_std_14)                                                               \
+  F(hip_std_17)                                                               \
+  F(hip_std_20)                                                               \
+  F(hip_std_23)

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()

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)

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()
+{
+}

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)

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)

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()
+{
+}

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)

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;
+}

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;
+}

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;
+}

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)

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;
+  }
+}

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")

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;
+}

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;
+}

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>
+)

Tests/HIP/WithDefs/inc_hip/inc_hip.h

@@ -0,0 +1 @@
+#define INC_HIP

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;
+}

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)

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)

Tests/RunCMake/CompilerLauncher/HIP-env-Build-stdout.txt

@@ -0,0 +1 @@
+.*-E env USED_LAUNCHER=1.*

Tests/RunCMake/CompilerLauncher/HIP-env-launch-Build-stdout.txt

@@ -0,0 +1 @@
+.*-E env USED_LAUNCHER=1.*

Tests/RunCMake/CompilerLauncher/HIP-env.cmake

@@ -0,0 +1 @@
+include(HIP-common.cmake)

Tests/RunCMake/CompilerLauncher/HIP-launch-Build-stdout.txt

@@ -0,0 +1 @@
+.*-E env USED_LAUNCHER=1.*

Tests/RunCMake/CompilerLauncher/HIP-launch-env.cmake

@@ -0,0 +1,3 @@
+set(CTEST_USE_LAUNCHERS 1)
+include(CTestUseLaunchers)
+include(HIP-env.cmake)

Tests/RunCMake/CompilerLauncher/HIP-launch.cmake

@@ -0,0 +1,3 @@
+set(CTEST_USE_LAUNCHERS 1)
+include(CTestUseLaunchers)
+include(HIP.cmake)

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)

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()

Tests/RunCMake/CompilerLauncher/main.hip

@@ -0,0 +1,4 @@
+int main()
+{
+  return 0;
+}