diff --git a/.gitlab-ci.yml b/.gitlab-ci.yml
index 236eac9..754b2a1 100644
--- a/.gitlab-ci.yml
+++ b/.gitlab-ci.yml
@@ -498,6 +498,31 @@
     variables:
         RSYNC_DESTINATION: dev
 
+## Sanitizer builds
+
+build:fedora34-asan:
+    extends:
+        - .fedora34_asan
+        - .cmake_build_linux
+        - .cmake_build_artifacts
+        - .linux_builder_tags_qt
+        - .run_manually
+    variables:
+        CMAKE_CI_JOB_NIGHTLY: "true"
+
+test:fedora34-asan:
+    extends:
+        - .fedora34_asan
+        - .cmake_memcheck_linux
+        - .linux_builder_tags_qt
+        - .run_automatically
+    dependencies:
+        - build:fedora34-asan
+    needs:
+        - build:fedora34-asan
+    variables:
+        CMAKE_CI_JOB_NIGHTLY: "true"
+
 # macOS builds
 
 build:macos-x86_64-ninja:
diff --git a/.gitlab/ci/configure_fedora34_asan.cmake b/.gitlab/ci/configure_fedora34_asan.cmake
new file mode 100644
index 0000000..c22cdb7
--- /dev/null
+++ b/.gitlab/ci/configure_fedora34_asan.cmake
@@ -0,0 +1,4 @@
+set(CMAKE_C_FLAGS "-fsanitize=address" CACHE STRING "")
+set(CMAKE_CXX_FLAGS "-fsanitize=address" CACHE STRING "")
+
+include("${CMAKE_CURRENT_LIST_DIR}/configure_fedora34_common.cmake")
diff --git a/.gitlab/ci/ctest_exclusions.cmake b/.gitlab/ci/ctest_exclusions.cmake
index b885a6a..3460d48 100644
--- a/.gitlab/ci/ctest_exclusions.cmake
+++ b/.gitlab/ci/ctest_exclusions.cmake
@@ -13,6 +13,13 @@
     "^ExternalProjectUpdateSetup$")
 endif ()
 
+if ("$ENV{CMAKE_CONFIGURATION}" MATCHES "_asan")
+  list(APPEND test_exclusions
+    CTestTest2 # crashes on purpose
+    BootstrapTest # no need to cover this for asan
+    )
+endif()
+
 string(REPLACE ";" "|" test_exclusions "${test_exclusions}")
 if (test_exclusions)
   set(test_exclusions "(${test_exclusions})")
diff --git a/.gitlab/ci/ctest_memcheck.cmake b/.gitlab/ci/ctest_memcheck.cmake
new file mode 100644
index 0000000..dac907c
--- /dev/null
+++ b/.gitlab/ci/ctest_memcheck.cmake
@@ -0,0 +1,45 @@
+cmake_minimum_required(VERSION 3.8)
+
+include("${CMAKE_CURRENT_LIST_DIR}/gitlab_ci.cmake")
+
+# Read the files from the build directory.
+ctest_read_custom_files("${CTEST_BINARY_DIRECTORY}")
+
+# Pick up from where the configure left off.
+ctest_start(APPEND)
+
+include(ProcessorCount)
+ProcessorCount(nproc)
+if (NOT "$ENV{CTEST_MAX_PARALLELISM}" STREQUAL "")
+  if (nproc GREATER "$ENV{CTEST_MAX_PARALLELISM}")
+    set(nproc "$ENV{CTEST_MAX_PARALLELISM}")
+  endif ()
+endif ()
+
+set(CTEST_MEMORYCHECK_TYPE "$ENV{CTEST_MEMORYCHECK_TYPE}")
+set(CTEST_MEMORYCHECK_SANITIZER_OPTIONS "$ENV{CTEST_MEMORYCHECK_SANITIZER_OPTIONS}")
+
+set(lsan_suppressions "${CMAKE_CURRENT_LIST_DIR}/ctest_memcheck_$ENV{CMAKE_CONFIGURATION}.lsan.supp")
+if (EXISTS "${lsan_suppressions}")
+  set(ENV{LSAN_OPTIONS} "suppressions='${lsan_suppressions}'")
+endif ()
+
+include("${CMAKE_CURRENT_LIST_DIR}/ctest_exclusions.cmake")
+ctest_memcheck(
+  PARALLEL_LEVEL "${nproc}"
+  TEST_LOAD "${nproc}"
+  RETURN_VALUE test_result
+  EXCLUDE "${test_exclusions}"
+  DEFECT_COUNT defects)
+ctest_submit(PARTS Test)
+ctest_submit(PARTS Memcheck)
+
+if (test_result)
+  message(FATAL_ERROR
+    "Failed to test")
+endif ()
+
+if (defects)
+  message(FATAL_ERROR
+    "Found ${defects} memcheck defects")
+endif ()
diff --git a/.gitlab/ci/ctest_memcheck_fedora34_asan.lsan.supp b/.gitlab/ci/ctest_memcheck_fedora34_asan.lsan.supp
new file mode 100644
index 0000000..8ec1a03
--- /dev/null
+++ b/.gitlab/ci/ctest_memcheck_fedora34_asan.lsan.supp
@@ -0,0 +1 @@
+# Add 'leak:<pattern>' lines here to suppress known leaks.
diff --git a/.gitlab/ci/docker/fedora34/install_deps.sh b/.gitlab/ci/docker/fedora34/install_deps.sh
index 7d099fe..146d51f 100755
--- a/.gitlab/ci/docker/fedora34/install_deps.sh
+++ b/.gitlab/ci/docker/fedora34/install_deps.sh
@@ -12,6 +12,7 @@
 # Install development tools.
 dnf install --setopt=install_weak_deps=False -y \
     clang-tools-extra \
+    compiler-rt \
     gcc-c++ \
     git-core \
     make
diff --git a/.gitlab/ci/env_fedora34_asan.sh b/.gitlab/ci/env_fedora34_asan.sh
new file mode 100644
index 0000000..e976486
--- /dev/null
+++ b/.gitlab/ci/env_fedora34_asan.sh
@@ -0,0 +1,2 @@
+export CC=/usr/bin/clang
+export CXX=/usr/bin/clang++
diff --git a/.gitlab/os-linux.yml b/.gitlab/os-linux.yml
index 015df4f..997beab 100644
--- a/.gitlab/os-linux.yml
+++ b/.gitlab/os-linux.yml
@@ -69,7 +69,7 @@
 ### Fedora
 
 .fedora34:
-    image: "kitware/cmake:ci-fedora34-x86_64-2021-06-03"
+    image: "kitware/cmake:ci-fedora34-x86_64-2021-06-07"
 
     variables:
         GIT_CLONE_PATH: "$CI_BUILDS_DIR/cmake ci/long file name for testing purposes"
@@ -158,6 +158,27 @@
         CMAKE_CONFIGURATION: debian10_makefiles_clang
         CMAKE_GENERATOR: "Unix Makefiles"
 
+### Sanitizers
+
+.fedora_memcheck:
+    variables:
+        CMAKE_BUILD_TYPE: RelWithDebInfo
+
+.fedora_asan_addon:
+    extends: .fedora_memcheck
+
+    variables:
+        CTEST_MEMORYCHECK_TYPE: AddressSanitizer
+        CTEST_MEMORYCHECK_SANITIZER_OPTIONS: ""
+
+.fedora34_asan:
+    extends:
+        - .fedora34
+        - .fedora_asan_addon
+
+    variables:
+        CMAKE_CONFIGURATION: fedora34_asan
+
 ### Intel Compiler
 
 .intelcompiler:
@@ -306,6 +327,15 @@
 
     interruptible: true
 
+.cmake_memcheck_linux:
+    stage: test
+
+    script:
+        - *before_script_linux
+        - "$LAUNCHER ctest --output-on-failure -V -S .gitlab/ci/ctest_memcheck.cmake"
+
+    interruptible: true
+
 .cmake_build_linux_release:
     stage: build
 
diff --git a/Help/command/ctest_test.rst b/Help/command/ctest_test.rst
index c61d01e..2153c90 100644
--- a/Help/command/ctest_test.rst
+++ b/Help/command/ctest_test.rst
@@ -248,5 +248,25 @@
 Attached Files
 """"""""""""""
 
-To associate other types of files with a test, use the
-:prop_test:`ATTACHED_FILES` or :prop_test:`ATTACHED_FILES_ON_FAIL` test properties.
+The following example demonstrates how to upload non-image files to CDash.
+
+.. code-block:: c++
+
+   std::cout <<
+     "<DartMeasurementFile type=\"file\" name=\"MyTestInputData\">" <<
+     "/dir/to/data.csv</DartMeasurementFile>" << std::endl;
+
+If the name of the file to upload is known at configure time, you can use the
+:prop_test:`ATTACHED_FILES` or :prop_test:`ATTACHED_FILES_ON_FAIL` test
+properties instead.
+
+Custom Details
+""""""""""""""
+
+The following example demonstrates how to specify a custom value for the
+``Test Details`` field displayed on CDash.
+
+.. code-block:: c++
+
+   std::cout <<
+     "<CTestDetails>My Custom Details Value</CTestDetails>" << std::endl;
diff --git a/Help/command/enable_language.rst b/Help/command/enable_language.rst
index ce765de..d2acbc8 100644
--- a/Help/command/enable_language.rst
+++ b/Help/command/enable_language.rst
@@ -10,7 +10,7 @@
 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 @@
 .. 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.
 
diff --git a/Help/command/project.rst b/Help/command/project.rst
index 8a6bc1e..2a9dcfe 100644
--- a/Help/command/project.rst
+++ b/Help/command/project.rst
@@ -103,7 +103,7 @@
 
   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.
diff --git a/Help/envvar/HIPCXX.rst b/Help/envvar/HIPCXX.rst
new file mode 100644
index 0000000..23329e9
--- /dev/null
+++ b/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.
diff --git a/Help/envvar/HIPFLAGS.rst b/Help/envvar/HIPFLAGS.rst
new file mode 100644
index 0000000..0df3416
--- /dev/null
+++ b/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>`.
diff --git a/Help/manual/cmake-env-variables.7.rst b/Help/manual/cmake-env-variables.7.rst
index bfdc841..fa38588 100644
--- a/Help/manual/cmake-env-variables.7.rst
+++ b/Help/manual/cmake-env-variables.7.rst
@@ -67,6 +67,8 @@
    /envvar/CXXFLAGS
    /envvar/FC
    /envvar/FFLAGS
+   /envvar/HIPCXX
+   /envvar/HIPFLAGS
    /envvar/ISPC
    /envvar/ISPCFLAGS
    /envvar/OBJC
diff --git a/Help/manual/cmake-generator-expressions.7.rst b/Help/manual/cmake-generator-expressions.7.rst
index 0ba41c8..05b7bc0 100644
--- a/Help/manual/cmake-generator-expressions.7.rst
+++ b/Help/manual/cmake-generator-expressions.7.rst
@@ -181,6 +181,13 @@
   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 @@
   ``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 @@
   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 @@
   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.
diff --git a/Help/manual/cmake-presets.7.rst b/Help/manual/cmake-presets.7.rst
index 0bcd691..db26b5a 100644
--- a/Help/manual/cmake-presets.7.rst
+++ b/Help/manual/cmake-presets.7.rst
@@ -401,6 +401,19 @@
   are applied. Setting a variable to ``null`` causes it to not be set,
   even if a value was inherited from another preset.
 
+  .. note::
+
+    For a CMake project using ExternalProject with a configuration preset
+    having environment variables needed in the ExternalProject, use a build
+    preset that inherits that configuration preset or the ExternalProject
+    will not have the environment variables set in the configuration preset.
+    Example: suppose the host defaults to one compiler (say Clang)
+    and the user wishes to use another compiler (say GCC). Set configuration
+    preset environment variables ``CC`` and ``CXX`` and use a build preset
+    that inherits that configuration preset. Otherwise the ExternalProject
+    may use a different (system default) compiler than the top-level CMake
+    project.
+
 ``configurePreset``
 
   An optional string specifying the name of a configure preset to
diff --git a/Help/manual/cmake-properties.7.rst b/Help/manual/cmake-properties.7.rst
index 0a3e36b..2db1ec2 100644
--- a/Help/manual/cmake-properties.7.rst
+++ b/Help/manual/cmake-properties.7.rst
@@ -212,6 +212,7 @@
    /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
diff --git a/Help/manual/cmake-variables.7.rst b/Help/manual/cmake-variables.7.rst
index 873bd89..39fbbed 100644
--- a/Help/manual/cmake-variables.7.rst
+++ b/Help/manual/cmake-variables.7.rst
@@ -523,6 +523,7 @@
    /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
diff --git a/Help/manual/ctest.1.rst b/Help/manual/ctest.1.rst
index 72ef259..ab819f9 100644
--- a/Help/manual/ctest.1.rst
+++ b/Help/manual/ctest.1.rst
@@ -322,11 +322,13 @@
  Set the interactive mode to ``0`` or ``1``.
 
  This option causes CTest to run tests in either an interactive mode
- or a non-interactive mode.  On Windows this means that in
- non-interactive mode, all system debug pop up windows are blocked.
- In dashboard mode (``Experimental``, ``Nightly``, ``Continuous``), the default
- is non-interactive.  When just running tests not for a dashboard the
- default is to allow popups and interactive debugging.
+ or a non-interactive mode.  In dashboard mode (``Experimental``, ``Nightly``,
+ ``Continuous``), the default is non-interactive.  In non-interactive mode,
+ the environment variable :envvar:`DASHBOARD_TEST_FROM_CTEST` is set.
+
+ Prior to CMake 3.11, interactive mode on Windows allowed system debug
+ popup windows to appear.  Now, due to CTest's use of ``libuv`` to launch
+ test processes, all system debug popup windows are always blocked.
 
 ``--no-label-summary``
  Disable timing summary information for labels.
diff --git a/Help/prop_sf/LANGUAGE.rst b/Help/prop_sf/LANGUAGE.rst
index f14c176..a9b5638 100644
--- a/Help/prop_sf/LANGUAGE.rst
+++ b/Help/prop_sf/LANGUAGE.rst
@@ -6,8 +6,8 @@
 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
diff --git a/Help/prop_tgt/HIP_ARCHITECTURES.rst b/Help/prop_tgt/HIP_ARCHITECTURES.rst
new file mode 100644
index 0000000..06f956b
--- /dev/null
+++ b/Help/prop_tgt/HIP_ARCHITECTURES.rst
@@ -0,0 +1,27 @@
+HIP_ARCHITECTURES
+-----------------
+
+.. versionadded:: 3.21
+
+List of AMD GPU architectures to generate device code for.
+
+A non-empty false value (e.g. ``OFF``) disables adding architectures.
+This is intended to support packagers and rare cases where full control
+over the passed flags is required.
+
+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``.
diff --git a/Help/prop_tgt/LANG_COMPILER_LAUNCHER.rst b/Help/prop_tgt/LANG_COMPILER_LAUNCHER.rst
index 16be3cd..cba8ac9 100644
--- a/Help/prop_tgt/LANG_COMPILER_LAUNCHER.rst
+++ b/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
diff --git a/Help/release/dev/ctest-measurements-docs.rst b/Help/release/dev/ctest-measurements-docs.rst
new file mode 100644
index 0000000..e47582e
--- /dev/null
+++ b/Help/release/dev/ctest-measurements-docs.rst
@@ -0,0 +1,5 @@
+ctest-measurements-docs
+-----------------------
+
+* :manual:`ctest(1)` gained documentation for its ability to capture
+  :ref:`Additional Test Measurements`.
diff --git a/Help/release/dev/ctest-runtime-files.rst b/Help/release/dev/ctest-runtime-files.rst
new file mode 100644
index 0000000..f13baa4
--- /dev/null
+++ b/Help/release/dev/ctest-runtime-files.rst
@@ -0,0 +1,8 @@
+ctest-runtime-files
+-------------------
+
+* :manual:`ctest(1)` learned to recognize files attached to a test at run time.
+  Previously it was only possible to attach files to tests at configure time
+  by using the :prop_test:`ATTACHED_FILES` or
+  :prop_test:`ATTACHED_FILES_ON_FAIL` test properties.
+  See :ref:`Additional Test Measurements` for more information.
diff --git a/Help/release/dev/hip.rst b/Help/release/dev/hip.rst
new file mode 100644
index 0000000..abf9a35
--- /dev/null
+++ b/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.
diff --git a/Help/variable/CMAKE_HIP_ARCHITECTURES.rst b/Help/variable/CMAKE_HIP_ARCHITECTURES.rst
new file mode 100644
index 0000000..0cf0201
--- /dev/null
+++ b/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 the default architecture chosen by the compiler.
+
+This variable is used to initialize the :prop_tgt:`HIP_ARCHITECTURES` property
+on all targets. See the target property for additional information.
diff --git a/Help/variable/CMAKE_LANG_COMPILER_ID.rst b/Help/variable/CMAKE_LANG_COMPILER_ID.rst
index 0abedde..f23b7a2 100644
--- a/Help/variable/CMAKE_LANG_COMPILER_ID.rst
+++ b/Help/variable/CMAKE_LANG_COMPILER_ID.rst
@@ -34,6 +34,7 @@
   OpenWatcom = Open Watcom (openwatcom.org)
   PGI = The Portland Group (pgroup.com)
   PathScale = PathScale (pathscale.com)
+  ROCMClang = ROCm Toolkit Clang-based Compiler (rocmdocs.amd.com)
   SDCC = Small Device C Compiler (sdcc.sourceforge.net)
   SunPro = Oracle Solaris Studio (oracle.com)
   TI = Texas Instruments (ti.com)
diff --git a/Help/variable/CMAKE_LANG_COMPILER_LAUNCHER.rst b/Help/variable/CMAKE_LANG_COMPILER_LAUNCHER.rst
index 89fd7bc..f16e594 100644
--- a/Help/variable/CMAKE_LANG_COMPILER_LAUNCHER.rst
+++ b/Help/variable/CMAKE_LANG_COMPILER_LAUNCHER.rst
@@ -6,7 +6,7 @@
 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.
diff --git a/Help/variable/CTEST_CUSTOM_MAXIMUM_FAILED_TEST_OUTPUT_SIZE.rst b/Help/variable/CTEST_CUSTOM_MAXIMUM_FAILED_TEST_OUTPUT_SIZE.rst
index 5aeae88..9f145c1 100644
--- a/Help/variable/CTEST_CUSTOM_MAXIMUM_FAILED_TEST_OUTPUT_SIZE.rst
+++ b/Help/variable/CTEST_CUSTOM_MAXIMUM_FAILED_TEST_OUTPUT_SIZE.rst
@@ -5,4 +5,7 @@
 will be collected by the :command:`ctest_test` command. Defaults to 307200
 (300 KiB).
 
+If a test's output contains the literal string "CTEST_FULL_OUTPUT",
+the output will not be truncated and may exceed the maximum size.
+
 .. include:: CTEST_CUSTOM_XXX.txt
diff --git a/Help/variable/CTEST_CUSTOM_MAXIMUM_PASSED_TEST_OUTPUT_SIZE.rst b/Help/variable/CTEST_CUSTOM_MAXIMUM_PASSED_TEST_OUTPUT_SIZE.rst
index 1fbb8c5..71ecf52 100644
--- a/Help/variable/CTEST_CUSTOM_MAXIMUM_PASSED_TEST_OUTPUT_SIZE.rst
+++ b/Help/variable/CTEST_CUSTOM_MAXIMUM_PASSED_TEST_OUTPUT_SIZE.rst
@@ -5,4 +5,7 @@
 will be collected by the :command:`ctest_test` command. Defaults to 1024
 (1 KiB).
 
+If a test's output contains the literal string "CTEST_FULL_OUTPUT",
+the output will not be truncated and may exceed the maximum size.
+
 .. include:: CTEST_CUSTOM_XXX.txt
diff --git a/Modules/CMakeCCompilerId.c.in b/Modules/CMakeCCompilerId.c.in
index 75e9d1a..1f19c00 100644
--- a/Modules/CMakeCCompilerId.c.in
+++ b/Modules/CMakeCCompilerId.c.in
@@ -11,6 +11,12 @@
 # define volatile
 #endif
 
+#if !defined(__has_include)
+/* If the compiler does not have __has_include, pretend the answer is
+   always no.  */
+#  define __has_include(x) 0
+#endif
+
 @CMAKE_C_COMPILER_ID_CONTENT@
 
 /* Construct the string literal in pieces to prevent the source from
diff --git a/Modules/CMakeCXXCompilerId.cpp.in b/Modules/CMakeCXXCompilerId.cpp.in
index a67caba..7362a08 100644
--- a/Modules/CMakeCXXCompilerId.cpp.in
+++ b/Modules/CMakeCXXCompilerId.cpp.in
@@ -5,6 +5,12 @@
 # error "A C compiler has been selected for C++."
 #endif
 
+#if !defined(__has_include)
+/* If the compiler does not have __has_include, pretend the answer is
+   always no.  */
+#  define __has_include(x) 0
+#endif
+
 @CMAKE_CXX_COMPILER_ID_CONTENT@
 
 /* Construct the string literal in pieces to prevent the source from
diff --git a/Modules/CMakeCompilerIdDetection.cmake b/Modules/CMakeCompilerIdDetection.cmake
index 850fc14..2197790 100644
--- a/Modules/CMakeCompilerIdDetection.cmake
+++ b/Modules/CMakeCompilerIdDetection.cmake
@@ -81,6 +81,11 @@
       ARMCC
       AppleClang
       ARMClang
+    )
+    if(NOT __skip_rocmclang)
+      list(APPEND ordered_compilers ROCMClang)
+    endif()
+    list(APPEND ordered_compilers
       Clang
       GNU
       MSVC
diff --git a/Modules/CMakeDetermineCompileFeatures.cmake b/Modules/CMakeDetermineCompileFeatures.cmake
index ee7fedb..a08e597 100644
--- a/Modules/CMakeDetermineCompileFeatures.cmake
+++ b/Modules/CMakeDetermineCompileFeatures.cmake
@@ -166,6 +166,62 @@
 
     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()
diff --git a/Modules/CMakeDetermineCompilerId.cmake b/Modules/CMakeDetermineCompilerId.cmake
index 448f071..bd1e732 100644
--- a/Modules/CMakeDetermineCompilerId.cmake
+++ b/Modules/CMakeDetermineCompilerId.cmake
@@ -150,6 +150,40 @@
     endif()
   endif()
 
+  # When invoked with HIPCC we need to extract the path to the underlying
+  # clang compiler when possible. This fixes the following issues:
+  #   env variables can change how hipcc behaves
+  #   allows us to properly find the binutils bundled with hip
+  if(CMAKE_${lang}_COMPILER_ID STREQUAL "ROCMClang"
+     AND CMAKE_${lang}_COMPILER MATCHES ".*hipcc")
+    get_filename_component(_hipcc_dir "${CMAKE_${lang}_COMPILER}" DIRECTORY)
+    execute_process(
+      COMMAND "${_hipcc_dir}/hipconfig"
+      --hipclangpath
+      OUTPUT_VARIABLE output
+      RESULT_VARIABLE result
+    )
+    if(result EQUAL 0 AND EXISTS "${output}")
+      if(lang STREQUAL "C")
+        set_property(CACHE CMAKE_${lang}_COMPILER PROPERTY VALUE "${output}/clang")
+        set(CMAKE_${lang}_COMPILER "${output}/clang" PARENT_SCOPE)
+      else()
+        set_property(CACHE CMAKE_${lang}_COMPILER PROPERTY VALUE "${output}/clang++")
+        set(CMAKE_${lang}_COMPILER "${output}/clang++" PARENT_SCOPE)
+      endif()
+    endif()
+    if(lang STREQUAL "HIP")
+      execute_process(
+        COMMAND "${_hipcc_dir}/hipconfig"
+        --rocmpath
+        OUTPUT_VARIABLE output
+        RESULT_VARIABLE result
+      )
+      if(result EQUAL 0)
+        set(_CMAKE_HIP_COMPILER_ROCM_ROOT "${output}" PARENT_SCOPE)
+      endif()
+    endif()
+  endif()
 
   if (COMPILER_QNXNTO AND CMAKE_${lang}_COMPILER_ID STREQUAL "GNU")
     execute_process(
diff --git a/Modules/CMakeDetermineHIPCompiler.cmake b/Modules/CMakeDetermineHIPCompiler.cmake
new file mode 100644
index 0000000..ed0110a
--- /dev/null
+++ b/Modules/CMakeDetermineHIPCompiler.cmake
@@ -0,0 +1,101 @@
+# 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)
+  # Analyze output from hipcc to get the current GPU architecture.
+  if(CMAKE_HIP_COMPILER_PRODUCED_OUTPUT MATCHES " -target-cpu ([a-z0-9]+) ")
+    set(CMAKE_HIP_ARCHITECTURES "${CMAKE_MATCH_1}" CACHE STRING "HIP architectures")
+  else()
+    message(FATAL_ERROR "Failed to find a working HIP architecture.")
+  endif()
+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")
diff --git a/Modules/CMakeHIPCompiler.cmake.in b/Modules/CMakeHIPCompiler.cmake.in
new file mode 100644
index 0000000..9a30a45
--- /dev/null
+++ b/Modules/CMakeHIPCompiler.cmake.in
@@ -0,0 +1,58 @@
+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 90)
+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_RUNTIME_LIBRARY_DEFAULT "SHARED")
+
+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@")
diff --git a/Modules/CMakeHIPCompilerABI.hip b/Modules/CMakeHIPCompilerABI.hip
new file mode 100644
index 0000000..6c912bd
--- /dev/null
+++ b/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;
+}
diff --git a/Modules/CMakeHIPCompilerId.hip.in b/Modules/CMakeHIPCompilerId.hip.in
new file mode 100644
index 0000000..5258efb
--- /dev/null
+++ b/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;
+}
diff --git a/Modules/CMakeHIPInformation.cmake b/Modules/CMakeHIPInformation.cmake
new file mode 100644
index 0000000..ec37e1c
--- /dev/null
+++ b/Modules/CMakeHIPInformation.cmake
@@ -0,0 +1,139 @@
+# 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)
+
+# Load the file and find the relevant HIP runtime.
+# This file will only exist after all compiler detection has finished
+include(${CMAKE_PLATFORM_INFO_DIR}/CMakeHIPRuntime.cmake OPTIONAL)
+if(COMMAND _CMAKE_FIND_HIP_RUNTIME)
+  _CMAKE_FIND_HIP_RUNTIME()
+endif()
diff --git a/Modules/CMakeHIPRuntime.cmake.in b/Modules/CMakeHIPRuntime.cmake.in
new file mode 100644
index 0000000..ade26bb
--- /dev/null
+++ b/Modules/CMakeHIPRuntime.cmake.in
@@ -0,0 +1,99 @@
+# Distributed under the OSI-approved BSD 3-Clause License.  See accompanying
+# file Copyright.txt or https://cmake.org/licensing for details.
+
+
+function(_CMAKE_FIND_HIP_RUNTIME )
+  # Determined when hipcc is the HIP compiler
+  set(_CMAKE_HIP_COMPILER_ROCM_ROOT "@_CMAKE_HIP_COMPILER_ROCM_ROOT@")
+
+  # Forward facing value that can be provided by the user
+  set(CMAKE_HIP_COMPILER_TOOLKIT_ROOT @CMAKE_HIP_COMPILER_TOOLKIT_ROOT@)
+
+  if(NOT DEFINED _CMAKE_HIP_DEVICE_RUNTIME_TARGET)
+    set(message_on_found TRUE)
+  endif()
+
+  set(explicit_search_only FALSE)
+  set(rocm_root_dirs )
+  if(DEFINED CMAKE_HIP_COMPILER_TOOLKIT_ROOT)
+    set(rocm_root_dirs "${CMAKE_HIP_COMPILER_TOOLKIT_ROOT}")
+    set(explicit_search_only TRUE)
+    set(error_message_location "the variable CMAKE_HIP_COMPILER_TOOLKIT_ROOT [\"${CMAKE_HIP_COMPILER_TOOLKIT_ROOT}\"]")
+  elseif(DEFINED ENV{CMAKE_HIP_COMPILER_TOOLKIT_ROOT})
+    set(rocm_root_dirs "$ENV{CMAKE_HIP_COMPILER_TOOLKIT_ROOT}")
+    set(explicit_search_only TRUE)
+    set(error_message_location "CMAKE_HIP_COMPILER_TOOLKIT_ROOT")
+    set(error_message_location "the environment variable CMAKE_HIP_COMPILER_TOOLKIT_ROOT [\"$ENV{CMAKE_HIP_COMPILER_TOOLKIT_ROOT}\"]")
+  elseif(DEFINED _CMAKE_HIP_COMPILER_ROCM_ROOT)
+    set(rocm_root_dirs "${_CMAKE_HIP_COMPILER_ROCM_ROOT}")
+    set(explicit_search_only TRUE)
+    set(error_message_location "the associated hipconfig --rocmpath [\"${_CMAKE_HIP_COMPILER_ROCM_ROOT}\"]")
+  endif()
+
+  # Guess on where rocm is installed
+  if(NOT rocm_root_dirs AND (UNIX AND NOT APPLE))
+    set(platform_base "/opt/rocm-")
+
+    # Finad all default rocm installations
+    file(GLOB possible_paths "${platform_base}*")
+
+    set(versions)
+    foreach(p ${possible_paths})
+      # Extract version number from end of string
+      string(REGEX MATCH "[0-9]+\\.[0-9]+\\.[0-9]+$" p_version ${p})
+      if(IS_DIRECTORY ${p} AND p_version)
+        list(APPEND versions ${p_version})
+      endif()
+    endforeach()
+
+    # Sort numerically in descending order, so we try the newest versions first.
+    list(SORT versions COMPARE NATURAL ORDER DESCENDING)
+
+    # With a descending list of versions, populate possible paths to search.
+    set(rocm_root_dirs "/opt/rocm")
+    foreach(v IN LISTS versions)
+      list(APPEND rocm_root_dirs "${platform_base}${v}")
+    endforeach()
+  endif()
+
+  set(search_rel_path "/lib/cmake/hip-lang/")
+  list(TRANSFORM rocm_root_dirs APPEND "${search_rel_path}")
+
+  find_package(hip-lang
+    CONFIG
+    PATHS ${rocm_root_dirs}
+    QUIET
+    NO_DEFAULT_PATH
+  )
+  if(NOT DEFINED _CMAKE_HIP_DEVICE_RUNTIME_TARGET AND NOT explicit_search_only)
+    find_package(hip-lang CONFIG QUIET)
+  endif()
+
+  if(DEFINED _CMAKE_HIP_DEVICE_RUNTIME_TARGET)
+    set(CMAKE_HIP_RUNTIME_LIBRARIES_STATIC
+      ${CMAKE_HIP_RUNTIME_LIBRARIES_STATIC}
+      ${_CMAKE_HIP_DEVICE_RUNTIME_TARGET} PARENT_SCOPE)
+    set(CMAKE_HIP_RUNTIME_LIBRARIES_SHARED
+      ${CMAKE_HIP_RUNTIME_LIBRARIES_SHARED}
+      ${_CMAKE_HIP_DEVICE_RUNTIME_TARGET} PARENT_SCOPE)
+  endif()
+
+  if(DEFINED _CMAKE_HIP_DEVICE_RUNTIME_TARGET AND message_on_found)
+    message(STATUS "Found HIP runtime: ${hip-lang_DIR}")
+  elseif(NOT DEFINED _CMAKE_HIP_DEVICE_RUNTIME_TARGET)
+    if(explicit_search_only)
+      set(error_message "Failed to find the HIP runtime, Could not find hip-lang-config.cmake at the following location(s):\n")
+      foreach(p IN LISTS rocm_root_dirs)
+        string(APPEND error_message "\t${p}\n")
+      endforeach()
+      string(APPEND "which are computed from the location specified by ${error_message_location}. \
+        Please specify CMAKE_HIP_COMPILER_TOOLKIT_ROOT to the location of")
+      message(FATAL_ERROR "${error_message}")
+    else()
+      message(FATAL_ERROR
+        "Failed to find the HIP runtime, Could not find hip-lang-config.cmake.\
+        Try setting CMAKE_HIP_COMPILER_TOOLKIT_ROOT")
+    endif()
+  endif()
+
+endfunction()
diff --git a/Modules/CMakeTestHIPCompiler.cmake b/Modules/CMakeTestHIPCompiler.cmake
new file mode 100644
index 0000000..d9fcc9d
--- /dev/null
+++ b/Modules/CMakeTestHIPCompiler.cmake
@@ -0,0 +1,119 @@
+# 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)
+
+
+# Setup the following:
+# Configure the new template file CMakeHipRuntime.cmake to
+# - ${CMAKE_PLATFORM_INFO_DIR}/
+# This file will do the actual find_package query. We than have
+# CMakeHIPInformation.cmake include `CMakeHipRuntime`
+# So it is included once system information has been finished
+#
+configure_file(
+ ${CMAKE_ROOT}/Modules/CMakeHIPRuntime.cmake.in
+ ${CMAKE_PLATFORM_INFO_DIR}/CMakeHIPRuntime.cmake
+ @ONLY
+)
+
+# 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)
+
+# Load the file and find the relevant HIP runtime.
+# This file will only exist after all compiler detection has finished
+include(${CMAKE_PLATFORM_INFO_DIR}/CMakeHIPRuntime.cmake)
+if(COMMAND _CMAKE_FIND_HIP_RUNTIME)
+  _CMAKE_FIND_HIP_RUNTIME()
+endif()
diff --git a/Modules/CheckLinkerFlag.cmake b/Modules/CheckLinkerFlag.cmake
index 3c7a828..28ac2e3 100644
--- a/Modules/CheckLinkerFlag.cmake
+++ b/Modules/CheckLinkerFlag.cmake
@@ -65,6 +65,8 @@
     set (_source "       program test\n       stop\n       end program")
   elseif (_lang MATCHES "CUDA")
     set (_source "__host__ int main() { return 0; }")
+  elseif (_lang MATCHES "HIP")
+    set (_source "__host__ int main() { return 0; }")
   elseif (_lang MATCHES "^(OBJC|OBJCXX)$")
     set (_source "#ifndef __OBJC__\n#  error \"Not an Objective-C++ compiler\"\n#endif\nint main(void) { return 0; }")
   else()
diff --git a/Modules/Compiler/CMakeCommonCompilerMacros.cmake b/Modules/Compiler/CMakeCommonCompilerMacros.cmake
index 29e6730..c86af98 100644
--- a/Modules/Compiler/CMakeCommonCompilerMacros.cmake
+++ b/Modules/Compiler/CMakeCommonCompilerMacros.cmake
@@ -170,3 +170,19 @@
     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()
diff --git a/Modules/Compiler/Clang-HIP.cmake b/Modules/Compiler/Clang-HIP.cmake
new file mode 100644
index 0000000..1030a43
--- /dev/null
+++ b/Modules/Compiler/Clang-HIP.cmake
@@ -0,0 +1,20 @@
+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  "")
+
+# Populated by CMakeHIPRuntime.cmake
+set(CMAKE_HIP_RUNTIME_LIBRARIES_STATIC "")
+set(CMAKE_HIP_RUNTIME_LIBRARIES_SHARED "")
diff --git a/Modules/Compiler/ROCMClang-ASM.cmake b/Modules/Compiler/ROCMClang-ASM.cmake
new file mode 100644
index 0000000..85d1110
--- /dev/null
+++ b/Modules/Compiler/ROCMClang-ASM.cmake
@@ -0,0 +1,2 @@
+include(Compiler/ROCMClang)
+__compiler_rocmclang(ASM)
diff --git a/Modules/Compiler/ROCMClang-C.cmake b/Modules/Compiler/ROCMClang-C.cmake
new file mode 100644
index 0000000..cdfa95d
--- /dev/null
+++ b/Modules/Compiler/ROCMClang-C.cmake
@@ -0,0 +1,7 @@
+include(Compiler/ROCMClang)
+__compiler_rocmclang(C)
+
+set(_rocm_clang_ver "${CMAKE_C_COMPILER_VERSION_INTERNAL}")
+set(CMAKE_C_COMPILER_VERSION "${CMAKE_C_COMPILER_VERSION_INTERNAL}")
+include(Compiler/Clang-C)
+set(CMAKE_C_COMPILER_VERSION "${_rocm_clang_ver}")
diff --git a/Modules/Compiler/ROCMClang-CXX.cmake b/Modules/Compiler/ROCMClang-CXX.cmake
new file mode 100644
index 0000000..5739c8e
--- /dev/null
+++ b/Modules/Compiler/ROCMClang-CXX.cmake
@@ -0,0 +1,7 @@
+include(Compiler/ROCMClang)
+__compiler_rocmclang(CXX)
+
+set(_rocm_clang_ver "${CMAKE_CXX_COMPILER_VERSION_INTERNAL}")
+set(CMAKE_CXX_COMPILER_VERSION "${CMAKE_CXX_COMPILER_VERSION_INTERNAL}")
+include(Compiler/Clang-CXX)
+set(CMAKE_CXX_COMPILER_VERSION "${_rocm_clang_ver}")
diff --git a/Modules/Compiler/ROCMClang-DetermineCompiler.cmake b/Modules/Compiler/ROCMClang-DetermineCompiler.cmake
new file mode 100644
index 0000000..c2fc99b
--- /dev/null
+++ b/Modules/Compiler/ROCMClang-DetermineCompiler.cmake
@@ -0,0 +1,19 @@
+
+set(_compiler_id_pp_test "defined(__clang__) && __has_include(<hip/hip_version.h>)")
+
+set(_compiler_id_version_compute "
+# if defined(__clang__) && __has_include(<hip/hip_version.h>)
+#  include <hip/hip_version.h>
+#  define @PREFIX@COMPILER_VERSION_MAJOR @MACRO_DEC@(HIP_VERSION_MAJOR)
+#  define @PREFIX@COMPILER_VERSION_MINOR @MACRO_DEC@(HIP_VERSION_MINOR)
+#  define @PREFIX@COMPILER_VERSION_PATCH @MACRO_DEC@(HIP_VERSION_PATCH)
+# endif")
+
+set(_compiler_id_simulate "
+# if defined(_MSC_VER)
+#  define @PREFIX@SIMULATE_ID \"MSVC\"
+# elif defined(__clang__)
+#  define @PREFIX@SIMULATE_ID \"Clang\"
+# elif defined(__GNUC__)
+#  define @PREFIX@SIMULATE_ID \"GNU\"
+# endif")
diff --git a/Modules/Compiler/ROCMClang-FindBinUtils.cmake b/Modules/Compiler/ROCMClang-FindBinUtils.cmake
new file mode 100644
index 0000000..e721c87
--- /dev/null
+++ b/Modules/Compiler/ROCMClang-FindBinUtils.cmake
@@ -0,0 +1 @@
+include(Compiler/Clang-FindBinUtils)
diff --git a/Modules/Compiler/ROCMClang-HIP.cmake b/Modules/Compiler/ROCMClang-HIP.cmake
new file mode 100644
index 0000000..7af7699
--- /dev/null
+++ b/Modules/Compiler/ROCMClang-HIP.cmake
@@ -0,0 +1,49 @@
+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  "")
+
+# Populated by CMakeHIPRuntime.cmake
+set(CMAKE_HIP_RUNTIME_LIBRARIES_STATIC "")
+set(CMAKE_HIP_RUNTIME_LIBRARIES_SHARED "")
+
+__compiler_check_default_language_standard(HIP 3.5 11)
diff --git a/Modules/Compiler/ROCMClang-OBJC.cmake b/Modules/Compiler/ROCMClang-OBJC.cmake
new file mode 100644
index 0000000..794973d
--- /dev/null
+++ b/Modules/Compiler/ROCMClang-OBJC.cmake
@@ -0,0 +1,7 @@
+include(Compiler/ROCMClang)
+__compiler_rocmclang(OBJC)
+
+set(_rocm_clang_ver "${CMAKE_OBJC_COMPILER_VERSION_INTERNAL}")
+set(CMAKE_OBJC_COMPILER_VERSION "${CMAKE_OBJC_COMPILER_VERSION_INTERNAL}")
+include(Compiler/Clang-OBJC)
+set(CMAKE_OBJC_COMPILER_VERSION "${_rocm_clang_ver}")
diff --git a/Modules/Compiler/ROCMClang-OBJCXX.cmake b/Modules/Compiler/ROCMClang-OBJCXX.cmake
new file mode 100644
index 0000000..82238e1
--- /dev/null
+++ b/Modules/Compiler/ROCMClang-OBJCXX.cmake
@@ -0,0 +1,7 @@
+include(Compiler/ROCMClang)
+__compiler_rocmclang(OBJCXX)
+
+set(_rocm_clang_ver "${CMAKE_OBJCXX_COMPILER_VERSION_INTERNAL}")
+set(CMAKE_OBJCXX_COMPILER_VERSION "${CMAKE_OBJCXX_COMPILER_VERSION_INTERNAL}")
+include(Compiler/Clang-OBJCXX)
+set(CMAKE_OBJCXX_COMPILER_VERSION "${_rocm_clang_ver}")
diff --git a/Modules/Compiler/ROCMClang.cmake b/Modules/Compiler/ROCMClang.cmake
new file mode 100644
index 0000000..6b38c2d
--- /dev/null
+++ b/Modules/Compiler/ROCMClang.cmake
@@ -0,0 +1,35 @@
+# Distributed under the OSI-approved BSD 3-Clause License.  See accompanying
+# file Copyright.txt or https://cmake.org/licensing for details.
+
+
+# This module is shared by multiple languages; use include blocker.
+include_guard()
+
+include(Compiler/CMakeCommonCompilerMacros)
+
+macro(__compiler_rocmclang lang)
+
+  set(CMAKE_${lang}_VERBOSE_FLAG "-v")
+
+  if(NOT "x${CMAKE_${lang}_SIMULATE_ID}" STREQUAL "xMSVC")
+    # Feature flags.
+    set(CMAKE_${lang}_COMPILE_OPTIONS_PIC "-fPIC")
+    set(CMAKE_${lang}_COMPILE_OPTIONS_PIE "-fPIE")
+    set(CMAKE_HIP_COMPILE_OPTIONS_VISIBILITY -fvisibility=)
+
+    string(APPEND CMAKE_HIP_FLAGS_INIT " ")
+    string(APPEND CMAKE_HIP_FLAGS_DEBUG_INIT " -g")
+    string(APPEND CMAKE_HIP_FLAGS_RELEASE_INIT " -O3 -DNDEBUG")
+    string(APPEND CMAKE_HIP_FLAGS_MINSIZEREL_INIT " -Os -DNDEBUG")
+    string(APPEND CMAKE_HIP_FLAGS_RELWITHDEBINFO_INIT " -O2 -g -DNDEBUG")
+  endif()
+
+  set(CMAKE_SHARED_LIBRARY_CREATE_HIP_FLAGS -shared)
+  set(CMAKE_INCLUDE_SYSTEM_FLAG_HIP "-isystem ")
+
+  set(CMAKE_${lang}_USE_RESPONSE_FILE_FOR_INCLUDES 1)
+  set(CMAKE_${lang}_USE_RESPONSE_FILE_FOR_LIBRARIES 1)
+  set(CMAKE_${lang}_USE_RESPONSE_FILE_FOR_OBJECTS 1)
+  set(CMAKE_${lang}_RESPONSE_FILE_FLAG "@")
+  set(CMAKE_${lang}_RESPONSE_FILE_LINK_FLAG "@")
+endmacro()
diff --git a/Modules/InstallRequiredSystemLibraries.cmake b/Modules/InstallRequiredSystemLibraries.cmake
index 2d08e08..fa7b125 100644
--- a/Modules/InstallRequiredSystemLibraries.cmake
+++ b/Modules/InstallRequiredSystemLibraries.cmake
@@ -80,7 +80,9 @@
         set(_Intel_archdir ia32)
       endif()
       set(_Intel_compiler_ver ${CMAKE_${LANG}_COMPILER_VERSION})
-      if(WIN32)
+      if(WIN32 AND EXISTS "${_Intel_basedir}/../redist/${_Intel_archdir}_win/compiler")
+        get_filename_component(_Intel_redistdir "${_Intel_basedir}/../redist/${_Intel_archdir}_win/compiler" ABSOLUTE)
+      elseif(WIN32)
         get_filename_component(_Intel_redistdir "${_Intel_basedir}/../../redist/${_Intel_archdir}/compiler" ABSOLUTE)
       elseif(APPLE)
         get_filename_component(_Intel_redistdir "${_Intel_basedir}/../../compiler/lib" ABSOLUTE)
diff --git a/Modules/Internal/CheckCompilerFlag.cmake b/Modules/Internal/CheckCompilerFlag.cmake
index 6b2a11e..9eb1bf0 100644
--- a/Modules/Internal/CheckCompilerFlag.cmake
+++ b/Modules/Internal/CheckCompilerFlag.cmake
@@ -24,6 +24,9 @@
   elseif(_lang STREQUAL "Fortran")
     set(_lang_src "       program test\n       stop\n       end program")
     set(_lang_fail_regex FAIL_REGEX "command[ -]line option .* is valid for .* but not for Fortran")
+  elseif(_lang STREQUAL "HIP")
+    set(_lang_src "__host__ int main() { return 0; }")
+    set(_lang_fail_regex FAIL_REGEX "argument unused during compilation: .*") # Clang
   elseif(_lang STREQUAL "OBJC")
     set(_lang_src [=[
 #ifndef __OBJC__
diff --git a/Modules/Internal/CheckSourceCompiles.cmake b/Modules/Internal/CheckSourceCompiles.cmake
index 3b2152a..8c3a418 100644
--- a/Modules/Internal/CheckSourceCompiles.cmake
+++ b/Modules/Internal/CheckSourceCompiles.cmake
@@ -22,6 +22,9 @@
     elseif(_lang STREQUAL "Fortran")
       set(_lang_textual "Fortran")
       set(_lang_ext "F90")
+    elseif(_lang STREQUAL "HIP")
+      set(_lang_textual "HIP")
+      set(_lang_ext "hip")
     elseif(_lang STREQUAL "ISPC")
       set(_lang_textual "ISPC")
       set(_lang_ext "ispc")
diff --git a/Modules/Internal/CheckSourceRuns.cmake b/Modules/Internal/CheckSourceRuns.cmake
index 676f3d0..75e9896 100644
--- a/Modules/Internal/CheckSourceRuns.cmake
+++ b/Modules/Internal/CheckSourceRuns.cmake
@@ -22,6 +22,9 @@
     elseif(_lang STREQUAL "Fortran")
       set(_lang_textual "Fortran")
       set(_lang_ext "F90")
+    elseif(_lang STREQUAL "HIP")
+      set(_lang_textual "HIP")
+      set(_lang_ext "hip")
     elseif(_lang STREQUAL "OBJC")
       set(_lang_textual "Objective-C")
       set(_lang_ext "m")
diff --git a/Modules/Internal/FeatureTesting.cmake b/Modules/Internal/FeatureTesting.cmake
index 72d96b3..b6f3c09 100644
--- a/Modules/Internal/FeatureTesting.cmake
+++ b/Modules/Internal/FeatureTesting.cmake
@@ -99,6 +99,16 @@
   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 @@
   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()
diff --git a/Modules/WriteCompilerDetectionHeader.cmake b/Modules/WriteCompilerDetectionHeader.cmake
index 0e4e028..54eb40e 100644
--- a/Modules/WriteCompilerDetectionHeader.cmake
+++ b/Modules/WriteCompilerDetectionHeader.cmake
@@ -384,6 +384,7 @@
     )
 endif()
 
+set(__skip_rocmclang TRUE)
 include(${CMAKE_CURRENT_LIST_DIR}/CMakeCompilerIdDetection.cmake)
 
 function(_load_compiler_variables CompilerId lang)
diff --git a/Source/CMakeVersion.cmake b/Source/CMakeVersion.cmake
index 7915a9c..8018e40 100644
--- a/Source/CMakeVersion.cmake
+++ b/Source/CMakeVersion.cmake
@@ -1,7 +1,7 @@
 # CMake version number components.
 set(CMake_VERSION_MAJOR 3)
 set(CMake_VERSION_MINOR 20)
-set(CMake_VERSION_PATCH 20210608)
+set(CMake_VERSION_PATCH 20210610)
 #set(CMake_VERSION_RC 0)
 set(CMake_VERSION_IS_DIRTY 0)
 
diff --git a/Source/CTest/cmCTestRunTest.cxx b/Source/CTest/cmCTestRunTest.cxx
index 5a6c775..a892113 100644
--- a/Source/CTest/cmCTestRunTest.cxx
+++ b/Source/CTest/cmCTestRunTest.cxx
@@ -40,6 +40,22 @@
 {
   cmCTestLog(this->CTest, HANDLER_VERBOSE_OUTPUT,
              this->GetIndex() << ": " << line << std::endl);
+
+  // Check for special CTest XML tags in this line of output.
+  // If any are found, this line is excluded from ProcessOutput.
+  if (!line.empty() && line.find("<CTest") != std::string::npos) {
+    if (this->TestHandler->CustomCompletionStatusRegex.find(line)) {
+      this->TestResult.CustomCompletionStatus =
+        this->TestHandler->CustomCompletionStatusRegex.match(1);
+      cmCTestLog(this->CTest, HANDLER_VERBOSE_OUTPUT,
+                 this->GetIndex() << ": "
+                                  << "Test Details changed to '"
+                                  << this->TestResult.CustomCompletionStatus
+                                  << "'" << std::endl);
+      return;
+    }
+  }
+
   this->ProcessOutput += line;
   this->ProcessOutput += "\n";
 
diff --git a/Source/CTest/cmCTestTestHandler.cxx b/Source/CTest/cmCTestTestHandler.cxx
index 1596d4a..730ec0f 100644
--- a/Source/CTest/cmCTestTestHandler.cxx
+++ b/Source/CTest/cmCTestTestHandler.cxx
@@ -308,6 +308,10 @@
   // regex to detect each individual <DartMeasurement>...</DartMeasurement>
   this->DartStuff1.compile(
     "(<DartMeasurement[^<]*</DartMeasurement[a-zA-Z]*>)");
+
+  // regex to detect <CTestDetails>...</CTestDetails>
+  this->CustomCompletionStatusRegex.compile(
+    "<CTestDetails>(.*)</CTestDetails>");
 }
 
 void cmCTestTestHandler::Initialize()
@@ -1460,7 +1464,11 @@
     xml.StartElement("NamedMeasurement");
     xml.Attribute("type", "text/string");
     xml.Attribute("name", "Completion Status");
-    xml.Element("Value", result.CompletionStatus);
+    if (result.CustomCompletionStatus.empty()) {
+      xml.Element("Value", result.CompletionStatus);
+    } else {
+      xml.Element("Value", result.CustomCompletionStatus);
+    }
     xml.EndElement(); // NamedMeasurement
 
     xml.StartElement("NamedMeasurement");
@@ -1550,19 +1558,29 @@
       result.Properties->AttachOnFail.end());
   }
   for (std::string const& file : result.Properties->AttachedFiles) {
-    const std::string& base64 = this->CTest->Base64GzipEncodeFile(file);
-    std::string const fname = cmSystemTools::GetFilenameName(file);
-    xml.StartElement("NamedMeasurement");
-    xml.Attribute("name", "Attached File");
-    xml.Attribute("encoding", "base64");
-    xml.Attribute("compression", "tar/gzip");
-    xml.Attribute("filename", fname);
-    xml.Attribute("type", "file");
-    xml.Element("Value", base64);
-    xml.EndElement(); // NamedMeasurement
+    this->AttachFile(xml, file, "");
   }
 }
 
+void cmCTestTestHandler::AttachFile(cmXMLWriter& xml, std::string const& file,
+                                    std::string const& name)
+{
+  const std::string& base64 = this->CTest->Base64GzipEncodeFile(file);
+  std::string const fname = cmSystemTools::GetFilenameName(file);
+  xml.StartElement("NamedMeasurement");
+  std::string measurement_name = name;
+  if (measurement_name.empty()) {
+    measurement_name = "Attached File";
+  }
+  xml.Attribute("name", measurement_name);
+  xml.Attribute("encoding", "base64");
+  xml.Attribute("compression", "tar/gzip");
+  xml.Attribute("filename", fname);
+  xml.Attribute("type", "file");
+  xml.Element("Value", base64);
+  xml.EndElement(); // NamedMeasurement
+}
+
 int cmCTestTestHandler::ExecuteCommands(std::vector<std::string>& vec)
 {
   for (std::string const& it : vec) {
@@ -2041,11 +2059,11 @@
         cmCTest::CleanString(measurementfile.match(5));
       if (cmSystemTools::FileExists(filename)) {
         long len = cmSystemTools::FileLength(filename);
+        std::string k1 = measurementfile.match(1);
+        std::string v1 = measurementfile.match(2);
+        std::string k2 = measurementfile.match(3);
+        std::string v2 = measurementfile.match(4);
         if (len == 0) {
-          std::string k1 = measurementfile.match(1);
-          std::string v1 = measurementfile.match(2);
-          std::string k2 = measurementfile.match(3);
-          std::string v2 = measurementfile.match(4);
           if (cmSystemTools::LowerCase(k1) == "type") {
             v1 = "text/string";
           }
@@ -2060,35 +2078,53 @@
           xml.Element("Value", "Image " + filename + " is empty");
           xml.EndElement();
         } else {
-          cmsys::ifstream ifs(filename.c_str(),
-                              std::ios::in
-#ifdef _WIN32
-                                | std::ios::binary
-#endif
-          );
-          auto file_buffer = cm::make_unique<unsigned char[]>(len + 1);
-          ifs.read(reinterpret_cast<char*>(file_buffer.get()), len);
-          auto encoded_buffer = cm::make_unique<unsigned char[]>(
-            static_cast<int>(static_cast<double>(len) * 1.5 + 5.0));
-
-          size_t rlen = cmsysBase64_Encode(file_buffer.get(), len,
-                                           encoded_buffer.get(), 1);
-
-          xml.StartElement("NamedMeasurement");
-          xml.Attribute(measurementfile.match(1).c_str(),
-                        measurementfile.match(2));
-          xml.Attribute(measurementfile.match(3).c_str(),
-                        measurementfile.match(4));
-          xml.Attribute("encoding", "base64");
-          std::ostringstream ostr;
-          for (size_t cc = 0; cc < rlen; cc++) {
-            ostr << encoded_buffer[cc];
-            if (cc % 60 == 0 && cc) {
-              ostr << std::endl;
-            }
+          std::string type;
+          std::string name;
+          if (cmSystemTools::LowerCase(k1) == "type") {
+            type = v1;
+          } else if (cmSystemTools::LowerCase(k2) == "type") {
+            type = v2;
           }
-          xml.Element("Value", ostr.str());
-          xml.EndElement(); // NamedMeasurement
+          if (cmSystemTools::LowerCase(k1) == "name") {
+            name = v1;
+          } else if (cmSystemTools::LowerCase(k2) == "name") {
+            name = v2;
+          }
+          if (type == "file") {
+            // Treat this measurement like an "ATTACHED_FILE" when the type
+            // is explicitly "file" (not an image).
+            this->AttachFile(xml, filename, name);
+          } else {
+            cmsys::ifstream ifs(filename.c_str(),
+                                std::ios::in
+#ifdef _WIN32
+                                  | std::ios::binary
+#endif
+            );
+            auto file_buffer = cm::make_unique<unsigned char[]>(len + 1);
+            ifs.read(reinterpret_cast<char*>(file_buffer.get()), len);
+            auto encoded_buffer = cm::make_unique<unsigned char[]>(
+              static_cast<int>(static_cast<double>(len) * 1.5 + 5.0));
+
+            size_t rlen = cmsysBase64_Encode(file_buffer.get(), len,
+                                             encoded_buffer.get(), 1);
+
+            xml.StartElement("NamedMeasurement");
+            xml.Attribute(measurementfile.match(1).c_str(),
+                          measurementfile.match(2));
+            xml.Attribute(measurementfile.match(3).c_str(),
+                          measurementfile.match(4));
+            xml.Attribute("encoding", "base64");
+            std::ostringstream ostr;
+            for (size_t cc = 0; cc < rlen; cc++) {
+              ostr << encoded_buffer[cc];
+              if (cc % 60 == 0 && cc) {
+                ostr << std::endl;
+              }
+            }
+            xml.Element("Value", ostr.str());
+            xml.EndElement(); // NamedMeasurement
+          }
         }
       } else {
         int idx = 4;
diff --git a/Source/CTest/cmCTestTestHandler.h b/Source/CTest/cmCTestTestHandler.h
index 6841624..bd51738 100644
--- a/Source/CTest/cmCTestTestHandler.h
+++ b/Source/CTest/cmCTestTestHandler.h
@@ -175,6 +175,7 @@
     std::string ExceptionStatus;
     bool CompressOutput;
     std::string CompletionStatus;
+    std::string CustomCompletionStatus;
     std::string Output;
     std::string DartString;
     int TestCount;
@@ -237,6 +238,8 @@
                              cmCTestTestResult const& result);
   // Write attached test files into the xml
   void AttachFiles(cmXMLWriter& xml, cmCTestTestResult& result);
+  void AttachFile(cmXMLWriter& xml, std::string const& file,
+                  std::string const& name);
 
   //! Clean test output to specified length
   void CleanTestOutput(std::string& output, size_t length);
@@ -356,6 +359,7 @@
   ListOfTests TestList;
   size_t TotalNumberOfTests;
   cmsys::RegularExpression DartStuff;
+  cmsys::RegularExpression CustomCompletionStatusRegex;
 
   std::ostream* LogFile;
 
diff --git a/Source/LexerParser/cmFortranParser.cxx b/Source/LexerParser/cmFortranParser.cxx
index 0ea3d97..3f3ddde 100644
--- a/Source/LexerParser/cmFortranParser.cxx
+++ b/Source/LexerParser/cmFortranParser.cxx
@@ -599,13 +599,13 @@
   /* YYRLINE[YYN] -- Source line where rule number YYN was defined.  */
 static const yytype_uint8 yyrline[] =
 {
-       0,    99,    99,    99,   102,   106,   111,   120,   126,   133,
-     138,   142,   147,   155,   160,   165,   170,   175,   180,   185,
-     190,   195,   199,   203,   207,   211,   212,   217,   217,   217,
-     218,   218,   219,   219,   220,   220,   221,   221,   222,   222,
-     223,   223,   224,   224,   225,   225,   226,   226,   229,   230,
-     231,   232,   233,   234,   235,   236,   237,   238,   239,   240,
-     241,   242,   243,   244,   245
+       0,   101,   101,   101,   104,   108,   113,   122,   128,   135,
+     140,   144,   149,   157,   162,   167,   172,   177,   182,   187,
+     192,   197,   201,   205,   209,   213,   214,   219,   219,   219,
+     220,   220,   221,   221,   222,   222,   223,   223,   224,   224,
+     225,   225,   226,   226,   227,   227,   228,   228,   231,   232,
+     233,   234,   235,   236,   237,   238,   239,   240,   241,   242,
+     243,   244,   245,   246,   247
 };
 #endif
 
@@ -1364,7 +1364,29 @@
   YY_SYMBOL_PRINT (yymsg, yykind, yyvaluep, yylocationp);
 
   YY_IGNORE_MAYBE_UNINITIALIZED_BEGIN
-  YY_USE (yykind);
+  switch (yykind)
+    {
+    case YYSYMBOL_STRING: /* STRING  */
+#line 95 "cmFortranParser.y"
+            { free(((*yyvaluep).string)); }
+#line 1373 "cmFortranParser.cxx"
+        break;
+
+    case YYSYMBOL_WORD: /* WORD  */
+#line 95 "cmFortranParser.y"
+            { free(((*yyvaluep).string)); }
+#line 1379 "cmFortranParser.cxx"
+        break;
+
+    case YYSYMBOL_CPP_INCLUDE_ANGLE: /* CPP_INCLUDE_ANGLE  */
+#line 95 "cmFortranParser.y"
+            { free(((*yyvaluep).string)); }
+#line 1385 "cmFortranParser.cxx"
+        break;
+
+      default:
+        break;
+    }
   YY_IGNORE_MAYBE_UNINITIALIZED_END
 }
 
@@ -1634,26 +1656,26 @@
   switch (yyn)
     {
   case 4: /* stmt: INTERFACE EOSTMT  */
-#line 102 "cmFortranParser.y"
+#line 104 "cmFortranParser.y"
                    {
     cmFortranParser* parser = cmFortran_yyget_extra(yyscanner);
     cmFortranParser_SetInInterface(parser, true);
   }
-#line 1643 "cmFortranParser.cxx"
+#line 1665 "cmFortranParser.cxx"
     break;
 
   case 5: /* stmt: USE WORD other EOSTMT  */
-#line 106 "cmFortranParser.y"
+#line 108 "cmFortranParser.y"
                         {
     cmFortranParser* parser = cmFortran_yyget_extra(yyscanner);
     cmFortranParser_RuleUse(parser, (yyvsp[-2].string));
     free((yyvsp[-2].string));
   }
-#line 1653 "cmFortranParser.cxx"
+#line 1675 "cmFortranParser.cxx"
     break;
 
   case 6: /* stmt: MODULE WORD other EOSTMT  */
-#line 111 "cmFortranParser.y"
+#line 113 "cmFortranParser.y"
                            {
     cmFortranParser* parser = cmFortran_yyget_extra(yyscanner);
     if (cmsysString_strcasecmp((yyvsp[-2].string), "function") != 0 &&
@@ -1663,22 +1685,22 @@
     }
     free((yyvsp[-2].string));
   }
-#line 1667 "cmFortranParser.cxx"
+#line 1689 "cmFortranParser.cxx"
     break;
 
   case 7: /* stmt: SUBMODULE LPAREN WORD RPAREN WORD other EOSTMT  */
-#line 120 "cmFortranParser.y"
+#line 122 "cmFortranParser.y"
                                                  {
     cmFortranParser* parser = cmFortran_yyget_extra(yyscanner);
     cmFortranParser_RuleSubmodule(parser, (yyvsp[-4].string), (yyvsp[-2].string));
     free((yyvsp[-4].string));
     free((yyvsp[-2].string));
   }
-#line 1678 "cmFortranParser.cxx"
+#line 1700 "cmFortranParser.cxx"
     break;
 
   case 8: /* stmt: SUBMODULE LPAREN WORD COLON WORD RPAREN WORD other EOSTMT  */
-#line 126 "cmFortranParser.y"
+#line 128 "cmFortranParser.y"
                                                             {
     cmFortranParser* parser = cmFortran_yyget_extra(yyscanner);
     cmFortranParser_RuleSubmoduleNested(parser, (yyvsp[-6].string), (yyvsp[-4].string), (yyvsp[-2].string));
@@ -1686,40 +1708,40 @@
     free((yyvsp[-4].string));
     free((yyvsp[-2].string));
   }
-#line 1690 "cmFortranParser.cxx"
+#line 1712 "cmFortranParser.cxx"
     break;
 
   case 9: /* stmt: INTERFACE WORD other EOSTMT  */
-#line 133 "cmFortranParser.y"
+#line 135 "cmFortranParser.y"
                               {
     cmFortranParser* parser = cmFortran_yyget_extra(yyscanner);
     cmFortranParser_SetInInterface(parser, true);
     free((yyvsp[-2].string));
   }
-#line 1700 "cmFortranParser.cxx"
+#line 1722 "cmFortranParser.cxx"
     break;
 
   case 10: /* stmt: END INTERFACE other EOSTMT  */
-#line 138 "cmFortranParser.y"
+#line 140 "cmFortranParser.y"
                              {
     cmFortranParser* parser = cmFortran_yyget_extra(yyscanner);
     cmFortranParser_SetInInterface(parser, false);
   }
-#line 1709 "cmFortranParser.cxx"
+#line 1731 "cmFortranParser.cxx"
     break;
 
   case 11: /* stmt: USE DCOLON WORD other EOSTMT  */
-#line 142 "cmFortranParser.y"
+#line 144 "cmFortranParser.y"
                                {
     cmFortranParser* parser = cmFortran_yyget_extra(yyscanner);
     cmFortranParser_RuleUse(parser, (yyvsp[-2].string));
     free((yyvsp[-2].string));
   }
-#line 1719 "cmFortranParser.cxx"
+#line 1741 "cmFortranParser.cxx"
     break;
 
   case 12: /* stmt: USE COMMA WORD DCOLON WORD other EOSTMT  */
-#line 147 "cmFortranParser.y"
+#line 149 "cmFortranParser.y"
                                           {
     if (cmsysString_strcasecmp((yyvsp[-4].string), "non_intrinsic") == 0) {
       cmFortranParser* parser = cmFortran_yyget_extra(yyscanner);
@@ -1728,139 +1750,139 @@
     free((yyvsp[-4].string));
     free((yyvsp[-2].string));
   }
-#line 1732 "cmFortranParser.cxx"
+#line 1754 "cmFortranParser.cxx"
     break;
 
   case 13: /* stmt: INCLUDE STRING other EOSTMT  */
-#line 155 "cmFortranParser.y"
+#line 157 "cmFortranParser.y"
                               {
     cmFortranParser* parser = cmFortran_yyget_extra(yyscanner);
     cmFortranParser_RuleInclude(parser, (yyvsp[-2].string));
     free((yyvsp[-2].string));
   }
-#line 1742 "cmFortranParser.cxx"
+#line 1764 "cmFortranParser.cxx"
     break;
 
   case 14: /* stmt: CPP_LINE_DIRECTIVE STRING other EOSTMT  */
-#line 160 "cmFortranParser.y"
+#line 162 "cmFortranParser.y"
                                          {
     cmFortranParser* parser = cmFortran_yyget_extra(yyscanner);
     cmFortranParser_RuleLineDirective(parser, (yyvsp[-2].string));
     free((yyvsp[-2].string));
   }
-#line 1752 "cmFortranParser.cxx"
+#line 1774 "cmFortranParser.cxx"
     break;
 
   case 15: /* stmt: CPP_INCLUDE_ANGLE other EOSTMT  */
-#line 165 "cmFortranParser.y"
+#line 167 "cmFortranParser.y"
                                  {
     cmFortranParser* parser = cmFortran_yyget_extra(yyscanner);
     cmFortranParser_RuleInclude(parser, (yyvsp[-2].string));
     free((yyvsp[-2].string));
   }
-#line 1762 "cmFortranParser.cxx"
+#line 1784 "cmFortranParser.cxx"
     break;
 
   case 16: /* stmt: include STRING other EOSTMT  */
-#line 170 "cmFortranParser.y"
+#line 172 "cmFortranParser.y"
                               {
     cmFortranParser* parser = cmFortran_yyget_extra(yyscanner);
     cmFortranParser_RuleInclude(parser, (yyvsp[-2].string));
     free((yyvsp[-2].string));
   }
-#line 1772 "cmFortranParser.cxx"
+#line 1794 "cmFortranParser.cxx"
     break;
 
   case 17: /* stmt: define WORD other EOSTMT  */
-#line 175 "cmFortranParser.y"
+#line 177 "cmFortranParser.y"
                            {
     cmFortranParser* parser = cmFortran_yyget_extra(yyscanner);
     cmFortranParser_RuleDefine(parser, (yyvsp[-2].string));
     free((yyvsp[-2].string));
   }
-#line 1782 "cmFortranParser.cxx"
+#line 1804 "cmFortranParser.cxx"
     break;
 
   case 18: /* stmt: undef WORD other EOSTMT  */
-#line 180 "cmFortranParser.y"
+#line 182 "cmFortranParser.y"
                           {
     cmFortranParser* parser = cmFortran_yyget_extra(yyscanner);
     cmFortranParser_RuleUndef(parser, (yyvsp[-2].string));
     free((yyvsp[-2].string));
   }
-#line 1792 "cmFortranParser.cxx"
+#line 1814 "cmFortranParser.cxx"
     break;
 
   case 19: /* stmt: ifdef WORD other EOSTMT  */
-#line 185 "cmFortranParser.y"
+#line 187 "cmFortranParser.y"
                           {
     cmFortranParser* parser = cmFortran_yyget_extra(yyscanner);
     cmFortranParser_RuleIfdef(parser, (yyvsp[-2].string));
     free((yyvsp[-2].string));
   }
-#line 1802 "cmFortranParser.cxx"
+#line 1824 "cmFortranParser.cxx"
     break;
 
   case 20: /* stmt: ifndef WORD other EOSTMT  */
-#line 190 "cmFortranParser.y"
+#line 192 "cmFortranParser.y"
                            {
     cmFortranParser* parser = cmFortran_yyget_extra(yyscanner);
     cmFortranParser_RuleIfndef(parser, (yyvsp[-2].string));
     free((yyvsp[-2].string));
   }
-#line 1812 "cmFortranParser.cxx"
+#line 1834 "cmFortranParser.cxx"
     break;
 
   case 21: /* stmt: if other EOSTMT  */
-#line 195 "cmFortranParser.y"
+#line 197 "cmFortranParser.y"
                   {
     cmFortranParser* parser = cmFortran_yyget_extra(yyscanner);
     cmFortranParser_RuleIf(parser);
   }
-#line 1821 "cmFortranParser.cxx"
+#line 1843 "cmFortranParser.cxx"
     break;
 
   case 22: /* stmt: elif other EOSTMT  */
-#line 199 "cmFortranParser.y"
+#line 201 "cmFortranParser.y"
                     {
     cmFortranParser* parser = cmFortran_yyget_extra(yyscanner);
     cmFortranParser_RuleElif(parser);
   }
-#line 1830 "cmFortranParser.cxx"
+#line 1852 "cmFortranParser.cxx"
     break;
 
   case 23: /* stmt: else other EOSTMT  */
-#line 203 "cmFortranParser.y"
+#line 205 "cmFortranParser.y"
                     {
     cmFortranParser* parser = cmFortran_yyget_extra(yyscanner);
     cmFortranParser_RuleElse(parser);
   }
-#line 1839 "cmFortranParser.cxx"
+#line 1861 "cmFortranParser.cxx"
     break;
 
   case 24: /* stmt: endif other EOSTMT  */
-#line 207 "cmFortranParser.y"
+#line 209 "cmFortranParser.y"
                      {
     cmFortranParser* parser = cmFortran_yyget_extra(yyscanner);
     cmFortranParser_RuleEndif(parser);
   }
-#line 1848 "cmFortranParser.cxx"
+#line 1870 "cmFortranParser.cxx"
     break;
 
   case 48: /* misc_code: WORD  */
-#line 229 "cmFortranParser.y"
+#line 231 "cmFortranParser.y"
                       { free ((yyvsp[0].string)); }
-#line 1854 "cmFortranParser.cxx"
+#line 1876 "cmFortranParser.cxx"
     break;
 
   case 55: /* misc_code: STRING  */
-#line 236 "cmFortranParser.y"
+#line 238 "cmFortranParser.y"
                       { free ((yyvsp[0].string)); }
-#line 1860 "cmFortranParser.cxx"
+#line 1882 "cmFortranParser.cxx"
     break;
 
 
-#line 1864 "cmFortranParser.cxx"
+#line 1886 "cmFortranParser.cxx"
 
       default: break;
     }
@@ -2085,6 +2107,6 @@
   return yyresult;
 }
 
-#line 248 "cmFortranParser.y"
+#line 250 "cmFortranParser.y"
 
 /* End of grammar */
diff --git a/Source/LexerParser/cmFortranParser.y b/Source/LexerParser/cmFortranParser.y
index e461160..a3e1c24 100644
--- a/Source/LexerParser/cmFortranParser.y
+++ b/Source/LexerParser/cmFortranParser.y
@@ -92,6 +92,8 @@
 %token SUBMODULE
 %token USE
 
+%destructor { free($$); } WORD STRING CPP_INCLUDE_ANGLE
+
 /*-------------------------------------------------------------------------*/
 /* grammar */
 %%
diff --git a/Source/cmCMakePresetsFileReadJSON.cxx b/Source/cmCMakePresetsFileReadJSON.cxx
index 909a78b..489551d 100644
--- a/Source/cmCMakePresetsFileReadJSON.cxx
+++ b/Source/cmCMakePresetsFileReadJSON.cxx
@@ -350,8 +350,8 @@
     ReadFileResult::READ_OK, ReadFileResult::INVALID_PRESET,
     PresetStringHelper);
 
-ReadFileResult PresetInheritsHelper(std::vector<std::string>& out,
-                                    const Json::Value* value)
+ReadFileResult PresetVectorOneOrMoreStringHelper(std::vector<std::string>& out,
+                                                 const Json::Value* value)
 {
   out.clear();
   if (!value) {
@@ -478,8 +478,8 @@
   cmJSONObjectHelper<ConfigurePreset, ReadFileResult>(
     ReadFileResult::READ_OK, ReadFileResult::INVALID_PRESET, false)
     .Bind("name"_s, &ConfigurePreset::Name, PresetStringHelper)
-    .Bind("inherits"_s, &ConfigurePreset::Inherits, PresetInheritsHelper,
-          false)
+    .Bind("inherits"_s, &ConfigurePreset::Inherits,
+          PresetVectorOneOrMoreStringHelper, false)
     .Bind("hidden"_s, &ConfigurePreset::Hidden, PresetBoolHelper, false)
     .Bind<std::nullptr_t>("vendor"_s, nullptr,
                           VendorHelper(ReadFileResult::INVALID_PRESET), false)
@@ -512,7 +512,8 @@
   cmJSONObjectHelper<BuildPreset, ReadFileResult>(
     ReadFileResult::READ_OK, ReadFileResult::INVALID_PRESET, false)
     .Bind("name"_s, &BuildPreset::Name, PresetStringHelper)
-    .Bind("inherits"_s, &BuildPreset::Inherits, PresetInheritsHelper, false)
+    .Bind("inherits"_s, &BuildPreset::Inherits,
+          PresetVectorOneOrMoreStringHelper, false)
     .Bind("hidden"_s, &BuildPreset::Hidden, PresetBoolHelper, false)
     .Bind<std::nullptr_t>("vendor"_s, nullptr,
                           VendorHelper(ReadFileResult::INVALID_PRESET), false)
@@ -528,7 +529,8 @@
           &BuildPreset::InheritConfigureEnvironment, PresetOptionalBoolHelper,
           false)
     .Bind("jobs"_s, &BuildPreset::Jobs, PresetOptionalIntHelper, false)
-    .Bind("targets"_s, &BuildPreset::Targets, PresetVectorStringHelper, false)
+    .Bind("targets"_s, &BuildPreset::Targets,
+          PresetVectorOneOrMoreStringHelper, false)
     .Bind("configuration"_s, &BuildPreset::Configuration, PresetStringHelper,
           false)
     .Bind("cleanFirst"_s, &BuildPreset::CleanFirst, PresetOptionalBoolHelper,
@@ -831,7 +833,8 @@
   cmJSONObjectHelper<TestPreset, ReadFileResult>(
     ReadFileResult::READ_OK, ReadFileResult::INVALID_PRESET, false)
     .Bind("name"_s, &TestPreset::Name, PresetStringHelper)
-    .Bind("inherits"_s, &TestPreset::Inherits, PresetInheritsHelper, false)
+    .Bind("inherits"_s, &TestPreset::Inherits,
+          PresetVectorOneOrMoreStringHelper, false)
     .Bind("hidden"_s, &TestPreset::Hidden, PresetBoolHelper, false)
     .Bind<std::nullptr_t>("vendor"_s, nullptr,
                           VendorHelper(ReadFileResult::INVALID_PRESET), false)
diff --git a/Source/cmCTest.cxx b/Source/cmCTest.cxx
index 7534e37..7c469c8 100644
--- a/Source/cmCTest.cxx
+++ b/Source/cmCTest.cxx
@@ -1612,8 +1612,33 @@
   return 0;
 }
 
+bool cmCTest::TryToChangeDirectory(std::string const& dir)
+{
+  cmCTestLog(this, OUTPUT,
+             "Internal ctest changing into directory: " << dir << std::endl);
+  cmsys::Status status = cmSystemTools::ChangeDirectory(dir);
+  if (!status) {
+    auto msg = "Failed to change working directory to \"" + dir +
+      "\" : " + status.GetString() + "\n";
+    cmCTestLog(this, ERROR_MESSAGE, msg);
+    return false;
+  }
+  return true;
+}
+
 std::string cmCTest::Base64GzipEncodeFile(std::string const& file)
 {
+  const std::string currDir = cmSystemTools::GetCurrentWorkingDirectory();
+  std::string parentDir = cmSystemTools::GetParentDirectory(file);
+
+  // Temporarily change to the file's directory so the tar gets created
+  // with a flat directory structure.
+  if (currDir != parentDir) {
+    if (!this->TryToChangeDirectory(parentDir)) {
+      return "";
+    }
+  }
+
   std::string tarFile = file + "_temp.tar.gz";
   std::vector<std::string> files;
   files.push_back(file);
@@ -1628,6 +1653,12 @@
   }
   std::string base64 = this->Base64EncodeFile(tarFile);
   cmSystemTools::RemoveFile(tarFile);
+
+  // Change back to the directory we started in.
+  if (currDir != parentDir) {
+    cmSystemTools::ChangeDirectory(currDir);
+  }
+
   return base64;
 }
 
@@ -2853,14 +2884,7 @@
     }
 
     if (currDir != workDir) {
-      cmCTestLog(this, OUTPUT,
-                 "Internal ctest changing into directory: " << workDir
-                                                            << std::endl);
-      cmsys::Status status = cmSystemTools::ChangeDirectory(workDir);
-      if (!status) {
-        auto msg = "Failed to change working directory to \"" + workDir +
-          "\" : " + status.GetString() + "\n";
-        cmCTestLog(this, ERROR_MESSAGE, msg);
+      if (!this->TryToChangeDirectory(workDir)) {
         return 1;
       }
     }
diff --git a/Source/cmCTest.h b/Source/cmCTest.h
index 392eb1c..551c116 100644
--- a/Source/cmCTest.h
+++ b/Source/cmCTest.h
@@ -536,6 +536,9 @@
   int RunCMakeAndTest(std::string* output);
   int ExecuteTests();
 
+  /** return true iff change directory was successful */
+  bool TryToChangeDirectory(std::string const& dir);
+
   struct Private;
   std::unique_ptr<Private> Impl;
 };
diff --git a/Source/cmComputeLinkDepends.cxx b/Source/cmComputeLinkDepends.cxx
index 4a6518fd..0b27e34 100644
--- a/Source/cmComputeLinkDepends.cxx
+++ b/Source/cmComputeLinkDepends.cxx
@@ -7,6 +7,7 @@
 #include <cstdio>
 #include <iterator>
 #include <sstream>
+#include <unordered_map>
 #include <utility>
 
 #include <cm/memory>
@@ -371,6 +372,12 @@
       // This target provides its own link interface information.
       this->AddLinkEntries(depender_index, iface->Libraries);
       this->AddLinkObjects(iface->Objects);
+      for (auto const& language : iface->Languages) {
+        auto runtimeEntries = iface->LanguageRuntimeLibraries.find(language);
+        if (runtimeEntries != iface->LanguageRuntimeLibraries.end()) {
+          this->AddLinkEntries(depender_index, runtimeEntries->second);
+        }
+      }
 
       if (isIface) {
         return;
@@ -516,6 +523,13 @@
     this->Target->GetLinkImplementation(this->Config);
   this->AddLinkEntries(-1, impl->Libraries);
   this->AddLinkObjects(impl->Objects);
+
+  for (auto const& language : impl->Languages) {
+    auto runtimeEntries = impl->LanguageRuntimeLibraries.find(language);
+    if (runtimeEntries != impl->LanguageRuntimeLibraries.end()) {
+      this->AddLinkEntries(-1, runtimeEntries->second);
+    }
+  }
   for (cmLinkItem const& wi : impl->WrongConfigLibraries) {
     this->CheckWrongConfigItem(wi);
   }
diff --git a/Source/cmComputeLinkInformation.cxx b/Source/cmComputeLinkInformation.cxx
index 2647998..d15da0c 100644
--- a/Source/cmComputeLinkInformation.cxx
+++ b/Source/cmComputeLinkInformation.cxx
@@ -585,10 +585,10 @@
     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);
     }
diff --git a/Source/cmCoreTryCompile.cxx b/Source/cmCoreTryCompile.cxx
index 5399fd0..8ecf264 100644
--- a/Source/cmCoreTryCompile.cxx
+++ b/Source/cmCoreTryCompile.cxx
@@ -189,6 +189,8 @@
 // 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 @@
 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 @@
   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 @@
     } 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 @@
     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 @@
       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 @@
       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 @@
         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 @@
     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 @@
     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 @@
                                      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 @@
     cState.WriteProperties(fout, targetName);
     cxxState.WriteProperties(fout, targetName);
     cudaState.WriteProperties(fout, targetName);
+    hipState.WriteProperties(fout, targetName);
     objcState.WriteProperties(fout, targetName);
     objcxxState.WriteProperties(fout, targetName);
 
diff --git a/Source/cmExtraCodeBlocksGenerator.cxx b/Source/cmExtraCodeBlocksGenerator.cxx
index f217201..df14261 100644
--- a/Source/cmExtraCodeBlocksGenerator.cxx
+++ b/Source/cmExtraCodeBlocksGenerator.cxx
@@ -366,10 +366,11 @@
               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);
             }
diff --git a/Source/cmGeneratorExpressionNode.cxx b/Source/cmGeneratorExpressionNode.cxx
index 8030b64..c608bf9 100644
--- a/Source/cmGeneratorExpressionNode.cxx
+++ b/Source/cmGeneratorExpressionNode.cxx
@@ -718,7 +718,7 @@
 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 @@
 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 @@
     { "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 @@
     { "OBJC_COMPILER_VERSION", &objcCompilerVersionNode },
     { "OBJCXX_COMPILER_VERSION", &objcxxCompilerVersionNode },
     { "Fortran_COMPILER_VERSION", &fortranCompilerVersionNode },
+    { "HIP_COMPILER_VERSION", &hipCompilerVersionNode },
     { "PLATFORM_ID", &platformIdNode },
     { "COMPILE_FEATURES", &compileFeaturesNode },
     { "CONFIGURATION", &configurationNode },
diff --git a/Source/cmGeneratorTarget.cxx b/Source/cmGeneratorTarget.cxx
index 5deb2df..f268c6c 100644
--- a/Source/cmGeneratorTarget.cxx
+++ b/Source/cmGeneratorTarget.cxx
@@ -697,6 +697,7 @@
   this->SourcesAreContextDependent = Tribool::Indeterminate;
   this->Objects.clear();
   this->VisitedConfigsForObjects.clear();
+  this->LinkImplMap.clear();
 }
 
 void cmGeneratorTarget::AddSourceCommon(const std::string& src, bool before)
@@ -979,7 +980,7 @@
     // 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);
     }
@@ -1120,7 +1121,7 @@
   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()) {
@@ -1234,6 +1235,20 @@
                               &dagChecker, result, excludeImported, language);
     }
 
+    cmLinkImplementation const* impl = this->GetLinkImplementation(config);
+    if (impl != nullptr) {
+      auto runtimeEntries = impl->LanguageRuntimeLibraries.find(language);
+      if (runtimeEntries != impl->LanguageRuntimeLibraries.end()) {
+        for (auto const& lib : runtimeEntries->second) {
+          if (lib.Target) {
+            handleSystemIncludesDep(this->LocalGenerator, lib.Target, config,
+                                    this, &dagChecker, result, excludeImported,
+                                    language);
+          }
+        }
+      }
+    }
+
     std::for_each(result.begin(), result.end(),
                   cmSystemTools::ConvertToUnixSlashes);
     std::sort(result.begin(), result.end());
@@ -1473,31 +1488,80 @@
   }
 }
 
+void addInterfaceEntry(cmGeneratorTarget const* headTarget,
+                       std::string const& config, std::string const& prop,
+                       std::string const& lang,
+                       cmGeneratorExpressionDAGChecker* dagChecker,
+                       EvaluatedTargetPropertyEntries& entries,
+                       bool usage_requirements_only,
+                       std::vector<cmLinkImplItem> const& libraries)
+{
+  for (cmLinkImplItem const& lib : libraries) {
+    if (lib.Target) {
+      EvaluatedTargetPropertyEntry ee(lib, lib.Backtrace);
+      // Pretend $<TARGET_PROPERTY:lib.Target,prop> appeared in our
+      // caller's property and hand-evaluate it as if it were compiled.
+      // Create a context as cmCompiledGeneratorExpression::Evaluate does.
+      cmGeneratorExpressionContext context(
+        headTarget->GetLocalGenerator(), config, false, headTarget, headTarget,
+        true, lib.Backtrace, lang);
+      cmExpandList(lib.Target->EvaluateInterfaceProperty(
+                     prop, &context, dagChecker, usage_requirements_only),
+                   ee.Values);
+      ee.ContextDependent = context.HadContextSensitiveCondition;
+      entries.Entries.emplace_back(std::move(ee));
+    }
+  }
+}
+
+// IncludeRuntimeInterface is used to break the cycle in computing
+// the necessary transitive dependencies of targets that can occur
+// now that we have implicit language runtime targets.
+//
+// To determine the set of languages that a target has we need to iterate
+// all the sources which includes transitive INTERFACE sources.
+// Therefore we can't determine what language runtimes are needed
+// for a target until after all sources are computed.
+//
+// Therefore while computing the applicable INTERFACE_SOURCES we
+// must ignore anything in LanguageRuntimeLibraries or we would
+// create a cycle ( INTERFACE_SOURCES requires LanguageRuntimeLibraries,
+// LanguageRuntimeLibraries requires INTERFACE_SOURCES).
+//
+enum class IncludeRuntimeInterface
+{
+  Yes,
+  No
+};
 void AddInterfaceEntries(cmGeneratorTarget const* headTarget,
                          std::string const& config, std::string const& prop,
                          std::string const& lang,
                          cmGeneratorExpressionDAGChecker* dagChecker,
                          EvaluatedTargetPropertyEntries& entries,
+                         IncludeRuntimeInterface searchRuntime,
                          bool usage_requirements_only = true)
 {
-  if (cmLinkImplementationLibraries const* impl =
-        headTarget->GetLinkImplementationLibraries(config)) {
-    entries.HadContextSensitiveCondition = impl->HadContextSensitiveCondition;
-    for (cmLinkImplItem const& lib : impl->Libraries) {
-      if (lib.Target) {
-        EvaluatedTargetPropertyEntry ee(lib, lib.Backtrace);
-        // Pretend $<TARGET_PROPERTY:lib.Target,prop> appeared in our
-        // caller's property and hand-evaluate it as if it were compiled.
-        // Create a context as cmCompiledGeneratorExpression::Evaluate does.
-        cmGeneratorExpressionContext context(
-          headTarget->GetLocalGenerator(), config, false, headTarget,
-          headTarget, true, lib.Backtrace, lang);
-        cmExpandList(lib.Target->EvaluateInterfaceProperty(
-                       prop, &context, dagChecker, usage_requirements_only),
-                     ee.Values);
-        ee.ContextDependent = context.HadContextSensitiveCondition;
-        entries.Entries.emplace_back(std::move(ee));
+  if (searchRuntime == IncludeRuntimeInterface::Yes) {
+    if (cmLinkImplementation const* impl =
+          headTarget->GetLinkImplementation(config)) {
+      entries.HadContextSensitiveCondition =
+        impl->HadContextSensitiveCondition;
+
+      auto runtimeLibIt = impl->LanguageRuntimeLibraries.find(lang);
+      if (runtimeLibIt != impl->LanguageRuntimeLibraries.end()) {
+        addInterfaceEntry(headTarget, config, prop, lang, dagChecker, entries,
+                          usage_requirements_only, runtimeLibIt->second);
       }
+      addInterfaceEntry(headTarget, config, prop, lang, dagChecker, entries,
+                        usage_requirements_only, impl->Libraries);
+    }
+  } else {
+    if (cmLinkImplementationLibraries const* impl =
+          headTarget->GetLinkImplementationLibraries(config)) {
+      entries.HadContextSensitiveCondition =
+        impl->HadContextSensitiveCondition;
+      addInterfaceEntry(headTarget, config, prop, lang, dagChecker, entries,
+                        usage_requirements_only, impl->Libraries);
     }
   }
 }
@@ -1655,7 +1719,8 @@
   // Collect INTERFACE_SOURCES of all direct link-dependencies.
   EvaluatedTargetPropertyEntries linkInterfaceSourcesEntries;
   AddInterfaceEntries(this, config, "INTERFACE_SOURCES", std::string(),
-                      &dagChecker, linkInterfaceSourcesEntries);
+                      &dagChecker, linkInterfaceSourcesEntries,
+                      IncludeRuntimeInterface::No, true);
   std::vector<std::string>::size_type numFilesBefore = files.size();
   bool contextDependentInterfaceSources = processSources(
     this, linkInterfaceSourcesEntries, files, uniqueSrcs, debugSources);
@@ -3350,6 +3415,29 @@
   }
 }
 
+void cmGeneratorTarget::AddHIPArchitectureFlags(std::string& flags) const
+{
+  const std::string& property = this->GetSafeProperty("HIP_ARCHITECTURES");
+
+  if (property.empty()) {
+    this->Makefile->IssueMessage(MessageType::FATAL_ERROR,
+                                 "HIP_ARCHITECTURES is empty for target \"" +
+                                   this->GetName() + "\".");
+  }
+
+  // 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 =
@@ -3569,7 +3657,7 @@
   }
 
   AddInterfaceEntries(this, config, "INTERFACE_INCLUDE_DIRECTORIES", lang,
-                      &dagChecker, entries);
+                      &dagChecker, entries, IncludeRuntimeInterface::Yes);
 
   if (this->Makefile->IsOn("APPLE")) {
     if (cmLinkImplementationLibraries const* impl =
@@ -3794,7 +3882,7 @@
     this, config, language, &dagChecker, this->CompileOptionsEntries);
 
   AddInterfaceEntries(this, config, "INTERFACE_COMPILE_OPTIONS", language,
-                      &dagChecker, entries);
+                      &dagChecker, entries, IncludeRuntimeInterface::Yes);
 
   processOptions(this, entries, result, uniqueOptions, debugOptions,
                  "compile options", OptionsParse::Shell);
@@ -3836,7 +3924,8 @@
     this, config, std::string(), &dagChecker, this->CompileFeaturesEntries);
 
   AddInterfaceEntries(this, config, "INTERFACE_COMPILE_FEATURES",
-                      std::string(), &dagChecker, entries);
+                      std::string(), &dagChecker, entries,
+                      IncludeRuntimeInterface::Yes);
 
   processOptions(this, entries, result, uniqueFeatures, debugFeatures,
                  "compile features", OptionsParse::None);
@@ -3880,7 +3969,7 @@
     this, config, language, &dagChecker, this->CompileDefinitionsEntries);
 
   AddInterfaceEntries(this, config, "INTERFACE_COMPILE_DEFINITIONS", language,
-                      &dagChecker, entries);
+                      &dagChecker, entries, IncludeRuntimeInterface::Yes);
 
   if (!config.empty()) {
     std::string configPropName =
@@ -3938,7 +4027,7 @@
     this, config, language, &dagChecker, this->PrecompileHeadersEntries);
 
   AddInterfaceEntries(this, config, "INTERFACE_PRECOMPILE_HEADERS", language,
-                      &dagChecker, entries);
+                      &dagChecker, entries, IncludeRuntimeInterface::Yes);
 
   std::vector<BT<std::string>> list;
   processOptions(this, entries, list, uniqueOptions, debugDefines,
@@ -4323,7 +4412,7 @@
     this, config, language, &dagChecker, this->LinkOptionsEntries);
 
   AddInterfaceEntries(this, config, "INTERFACE_LINK_OPTIONS", language,
-                      &dagChecker, entries,
+                      &dagChecker, entries, IncludeRuntimeInterface::Yes,
                       this->GetPolicyStatusCMP0099() != cmPolicies::NEW);
 
   processOptions(this, entries, result, uniqueOptions, debugOptions,
@@ -4582,7 +4671,7 @@
     this, config, language, &dagChecker, this->LinkDirectoriesEntries);
 
   AddInterfaceEntries(this, config, "INTERFACE_LINK_DIRECTORIES", language,
-                      &dagChecker, entries,
+                      &dagChecker, entries, IncludeRuntimeInterface::Yes,
                       this->GetPolicyStatusCMP0099() != cmPolicies::NEW);
 
   processLinkDirectories(this, entries, result, uniqueDirectories,
@@ -4621,7 +4710,7 @@
     }
   }
   AddInterfaceEntries(this, config, "INTERFACE_LINK_DEPENDS", language,
-                      &dagChecker, entries,
+                      &dagChecker, entries, IncludeRuntimeInterface::Yes,
                       this->GetPolicyStatusCMP0099() != cmPolicies::NEW);
 
   processOptions(this, entries, result, uniqueOptions, false, "link depends",
@@ -4741,7 +4830,8 @@
       }
 
       // 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";
         }
@@ -6363,6 +6453,7 @@
     iface.AllDone = true;
     if (iface.Exists) {
       this->ComputeLinkInterface(config, iface, head, secondPass);
+      this->ComputeLinkInterfaceRuntimeLibraries(config, iface);
     }
   }
 
@@ -6886,6 +6977,83 @@
   }
 }
 
+namespace {
+
+template <typename ReturnType>
+ReturnType constructItem(cmGeneratorTarget* target,
+                         cmListFileBacktrace const& bt);
+
+template <>
+inline cmLinkImplItem constructItem(cmGeneratorTarget* target,
+                                    cmListFileBacktrace const& bt)
+{
+  return cmLinkImplItem(cmLinkItem(target, false, bt), false);
+}
+
+template <>
+inline cmLinkItem constructItem(cmGeneratorTarget* target,
+                                cmListFileBacktrace const& bt)
+{
+  return cmLinkItem(target, false, bt);
+}
+
+template <typename ValueType>
+std::vector<ValueType> computeImplicitLanguageTargets(
+  std::string const& lang, std::string const& config,
+  cmGeneratorTarget const* currentTarget)
+{
+  cmListFileBacktrace bt;
+  std::vector<ValueType> result;
+  cmLocalGenerator* lg = currentTarget->GetLocalGenerator();
+
+  std::string const& runtimeLibrary =
+    currentTarget->GetRuntimeLinkLibrary(lang, config);
+  if (cmProp runtimeLinkOptions = currentTarget->Makefile->GetDefinition(
+        "CMAKE_" + lang + "_RUNTIME_LIBRARIES_" + runtimeLibrary)) {
+    std::vector<std::string> libsVec = cmExpandedList(*runtimeLinkOptions);
+    result.reserve(libsVec.size());
+
+    for (std::string const& i : libsVec) {
+      cmGeneratorTarget::TargetOrString resolved =
+        currentTarget->ResolveTargetReference(i, lg);
+      if (resolved.Target) {
+        result.emplace_back(constructItem<ValueType>(resolved.Target, bt));
+      }
+    }
+  }
+
+  return result;
+}
+}
+
+void cmGeneratorTarget::ComputeLinkInterfaceRuntimeLibraries(
+  const std::string& config, cmOptionalLinkInterface& iface) const
+{
+  for (std::string const& lang : iface.Languages) {
+    if ((lang == "CUDA" || lang == "HIP") &&
+        iface.LanguageRuntimeLibraries.find(lang) ==
+          iface.LanguageRuntimeLibraries.end()) {
+      auto implicitTargets =
+        computeImplicitLanguageTargets<cmLinkItem>(lang, config, this);
+      iface.LanguageRuntimeLibraries[lang] = std::move(implicitTargets);
+    }
+  }
+}
+
+void cmGeneratorTarget::ComputeLinkImplementationRuntimeLibraries(
+  const std::string& config, cmOptionalLinkImplementation& impl) const
+{
+  for (std::string const& lang : impl.Languages) {
+    if ((lang == "CUDA" || lang == "HIP") &&
+        impl.LanguageRuntimeLibraries.find(lang) ==
+          impl.LanguageRuntimeLibraries.end()) {
+      auto implicitTargets =
+        computeImplicitLanguageTargets<cmLinkImplItem>(lang, config, this);
+      impl.LanguageRuntimeLibraries[lang] = std::move(implicitTargets);
+    }
+  }
+}
+
 const cmLinkInterface* cmGeneratorTarget::GetImportLinkInterface(
   const std::string& config, cmGeneratorTarget const* headTarget,
   bool usage_requirements_only, bool secondPass) const
@@ -7150,6 +7318,7 @@
   if (!impl.LanguagesDone) {
     impl.LanguagesDone = true;
     this->ComputeLinkImplementationLanguages(config, impl);
+    this->ComputeLinkImplementationRuntimeLibraries(config, impl);
   }
   return &impl;
 }
diff --git a/Source/cmGeneratorTarget.h b/Source/cmGeneratorTarget.h
index be36f3f..ed66fb1 100644
--- a/Source/cmGeneratorTarget.h
+++ b/Source/cmGeneratorTarget.h
@@ -459,6 +459,8 @@
   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(
@@ -1109,6 +1111,12 @@
   cmProp GetPropertyWithPairedLanguageSupport(std::string const& lang,
                                               const char* suffix) const;
 
+  void ComputeLinkImplementationRuntimeLibraries(
+    const std::string& config, cmOptionalLinkImplementation& impl) const;
+
+  void ComputeLinkInterfaceRuntimeLibraries(
+    const std::string& config, cmOptionalLinkInterface& iface) const;
+
 public:
   const std::vector<const cmGeneratorTarget*>& GetLinkImplementationClosure(
     const std::string& config) const;
diff --git a/Source/cmLinkItem.h b/Source/cmLinkItem.h
index db44938..0863edd 100644
--- a/Source/cmLinkItem.h
+++ b/Source/cmLinkItem.h
@@ -7,6 +7,7 @@
 #include <map>
 #include <ostream>
 #include <string>
+#include <unordered_map>
 #include <vector>
 
 #include <cmext/algorithm>
@@ -80,6 +81,8 @@
 {
   // Languages whose runtime libraries must be linked.
   std::vector<std::string> Languages;
+  std::unordered_map<std::string, std::vector<cmLinkItem>>
+    LanguageRuntimeLibraries;
 
   // Shared library dependencies needed for linking on some platforms.
   std::vector<cmLinkItem> SharedDeps;
@@ -115,6 +118,8 @@
 {
   // Languages whose runtime libraries must be linked.
   std::vector<std::string> Languages;
+  std::unordered_map<std::string, std::vector<cmLinkImplItem>>
+    LanguageRuntimeLibraries;
 
   // Whether the list depends on a link language genex.
   bool HadLinkLanguageSensitiveCondition = false;
diff --git a/Source/cmLocalGenerator.cxx b/Source/cmLocalGenerator.cxx
index 6a49b84..3b282de 100644
--- a/Source/cmLocalGenerator.cxx
+++ b/Source/cmLocalGenerator.cxx
@@ -780,9 +780,9 @@
     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 @@
     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 @@
       compilerSimulateId =
         this->Makefile->GetSafeDefinition("CMAKE_CXX_SIMULATE_ID");
     }
+  } else if (lang == "HIP") {
+    target->AddHIPArchitectureFlags(flags);
   }
 
   // Add VFS Overlay for Clang compilers
diff --git a/Source/cmLocalUnixMakefileGenerator3.cxx b/Source/cmLocalUnixMakefileGenerator3.cxx
index 3a65a80..0667c55 100644
--- a/Source/cmLocalUnixMakefileGenerator3.cxx
+++ b/Source/cmLocalUnixMakefileGenerator3.cxx
@@ -289,9 +289,9 @@
     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 @@
     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);
     }
diff --git a/Source/cmMakefileTargetGenerator.cxx b/Source/cmMakefileTargetGenerator.cxx
index e87f81b..6d8376c 100644
--- a/Source/cmMakefileTargetGenerator.cxx
+++ b/Source/cmMakefileTargetGenerator.cxx
@@ -945,7 +945,8 @@
     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)) {
diff --git a/Source/cmNinjaTargetGenerator.cxx b/Source/cmNinjaTargetGenerator.cxx
index 4feb0bb..6080270 100644
--- a/Source/cmNinjaTargetGenerator.cxx
+++ b/Source/cmNinjaTargetGenerator.cxx
@@ -833,7 +833,8 @@
   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)) {
diff --git a/Source/cmStandardLevelResolver.cxx b/Source/cmStandardLevelResolver.cxx
index b198a00..73d0fed 100644
--- a/Source/cmStandardLevelResolver.cxx
+++ b/Source/cmStandardLevelResolver.cxx
@@ -36,6 +36,9 @@
 
 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 @@
 };
 
 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 @@
       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 @@
     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";
diff --git a/Source/cmTarget.cxx b/Source/cmTarget.cxx
index 6f2e447..7622700 100644
--- a/Source/cmTarget.cxx
+++ b/Source/cmTarget.cxx
@@ -288,6 +288,7 @@
     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 @@
     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 @@
   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 @@
     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 @@
     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 {
diff --git a/Source/cmake.cxx b/Source/cmake.cxx
index 9a866ab..7868859 100644
--- a/Source/cmake.cxx
+++ b/Source/cmake.cxx
@@ -218,6 +218,7 @@
     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 @@
   // 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;
@@ -3273,7 +3276,9 @@
     this->UnprocessedPresetEnvironment = expandedPreset->Environment;
     this->ProcessPresetEnvironment();
 
-    if (jobs == cmake::DEFAULT_BUILD_PARALLEL_LEVEL && expandedPreset->Jobs) {
+    if ((jobs == cmake::DEFAULT_BUILD_PARALLEL_LEVEL ||
+         jobs == cmake::NO_BUILD_PARALLEL_LEVEL) &&
+        expandedPreset->Jobs) {
       jobs = *expandedPreset->Jobs;
     }
 
diff --git a/Source/cmake.h b/Source/cmake.h
index e845662..5a2a88f 100644
--- a/Source/cmake.h
+++ b/Source/cmake.h
@@ -297,7 +297,7 @@
     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 @@
   FileExtensions CudaFileExtensions;
   FileExtensions ISPCFileExtensions;
   FileExtensions FortranFileExtensions;
+  FileExtensions HipFileExtensions;
   bool ClearBuildSystem = false;
   bool DebugTryCompile = false;
   bool RegenerateDuringBuild = false;
@@ -838,3 +839,11 @@
   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)
diff --git a/Templates/MSBuild/FlagTables/v10_Link.json b/Templates/MSBuild/FlagTables/v10_Link.json
index ac5b8b1..494774b 100644
--- a/Templates/MSBuild/FlagTables/v10_Link.json
+++ b/Templates/MSBuild/FlagTables/v10_Link.json
@@ -1090,15 +1090,6 @@
     ]
   },
   {
-    "name": "MergeSections",
-    "switch": "MERGE:",
-    "comment": "Merge Sections",
-    "value": "",
-    "flags": [
-      "UserValue"
-    ]
-  },
-  {
     "name": "LinkKeyFile",
     "switch": "KEYFILE:",
     "comment": "Key File",
diff --git a/Templates/MSBuild/FlagTables/v11_Link.json b/Templates/MSBuild/FlagTables/v11_Link.json
index 5d5c13f..0db5d9e 100644
--- a/Templates/MSBuild/FlagTables/v11_Link.json
+++ b/Templates/MSBuild/FlagTables/v11_Link.json
@@ -1225,15 +1225,6 @@
     ]
   },
   {
-    "name": "MergeSections",
-    "switch": "MERGE:",
-    "comment": "Merge Sections",
-    "value": "",
-    "flags": [
-      "UserValue"
-    ]
-  },
-  {
     "name": "LinkKeyFile",
     "switch": "KEYFILE:",
     "comment": "Key File",
diff --git a/Templates/MSBuild/FlagTables/v12_Link.json b/Templates/MSBuild/FlagTables/v12_Link.json
index 5d5c13f..0db5d9e 100644
--- a/Templates/MSBuild/FlagTables/v12_Link.json
+++ b/Templates/MSBuild/FlagTables/v12_Link.json
@@ -1225,15 +1225,6 @@
     ]
   },
   {
-    "name": "MergeSections",
-    "switch": "MERGE:",
-    "comment": "Merge Sections",
-    "value": "",
-    "flags": [
-      "UserValue"
-    ]
-  },
-  {
     "name": "LinkKeyFile",
     "switch": "KEYFILE:",
     "comment": "Key File",
diff --git a/Templates/MSBuild/FlagTables/v140_Link.json b/Templates/MSBuild/FlagTables/v140_Link.json
index 3fb072c..7e293b1 100644
--- a/Templates/MSBuild/FlagTables/v140_Link.json
+++ b/Templates/MSBuild/FlagTables/v140_Link.json
@@ -1269,15 +1269,6 @@
     ]
   },
   {
-    "name": "MergeSections",
-    "switch": "MERGE:",
-    "comment": "Merge Sections",
-    "value": "",
-    "flags": [
-      "UserValue"
-    ]
-  },
-  {
     "name": "LinkKeyFile",
     "switch": "KEYFILE:",
     "comment": "Key File",
diff --git a/Templates/MSBuild/FlagTables/v141_Link.json b/Templates/MSBuild/FlagTables/v141_Link.json
index 66ee76f..e311240 100644
--- a/Templates/MSBuild/FlagTables/v141_Link.json
+++ b/Templates/MSBuild/FlagTables/v141_Link.json
@@ -1276,15 +1276,6 @@
     ]
   },
   {
-    "name": "MergeSections",
-    "switch": "MERGE:",
-    "comment": "Merge Sections",
-    "value": "",
-    "flags": [
-      "UserValue"
-    ]
-  },
-  {
     "name": "LinkKeyFile",
     "switch": "KEYFILE:",
     "comment": "Key File",
diff --git a/Templates/MSBuild/FlagTables/v142_Link.json b/Templates/MSBuild/FlagTables/v142_Link.json
index 110dcc2..d5fc2e4 100644
--- a/Templates/MSBuild/FlagTables/v142_Link.json
+++ b/Templates/MSBuild/FlagTables/v142_Link.json
@@ -1283,15 +1283,6 @@
     ]
   },
   {
-    "name": "MergeSections",
-    "switch": "MERGE:",
-    "comment": "Merge Sections",
-    "value": "",
-    "flags": [
-      "UserValue"
-    ]
-  },
-  {
     "name": "LinkKeyFile",
     "switch": "KEYFILE:",
     "comment": "Key File",
diff --git a/Tests/CMakeLists.txt b/Tests/CMakeLists.txt
index fccf15b..57d75e1 100644
--- a/Tests/CMakeLists.txt
+++ b/Tests/CMakeLists.txt
@@ -1485,6 +1485,10 @@
     add_subdirectory(CudaOnly)
   endif()
 
+  if(CMake_TEST_HIP)
+    add_subdirectory(HIP)
+  endif()
+
   if(CMake_TEST_ISPC)
     add_subdirectory(ISPC)
   endif()
diff --git a/Tests/HIP/ArchitectureOff/CMakeLists.txt b/Tests/HIP/ArchitectureOff/CMakeLists.txt
new file mode 100644
index 0000000..bccb3b4
--- /dev/null
+++ b/Tests/HIP/ArchitectureOff/CMakeLists.txt
@@ -0,0 +1,8 @@
+cmake_minimum_required(VERSION 3.18)
+project(HIPArchitecture HIP)
+
+# Make sure CMake doesn't pass architectures if HIP_ARCHITECTURES is OFF.
+string(APPEND CMAKE_HIP_FLAGS " --offload-arch=gfx908")
+
+add_executable(HIPOnlyArchitectureOff main.hip)
+set_property(TARGET HIPOnlyArchitectureOff PROPERTY HIP_ARCHITECTURES OFF)
diff --git a/Tests/HIP/ArchitectureOff/main.hip b/Tests/HIP/ArchitectureOff/main.hip
new file mode 100644
index 0000000..9256318
--- /dev/null
+++ b/Tests/HIP/ArchitectureOff/main.hip
@@ -0,0 +1,9 @@
+#ifdef __HIP_DEVICE_COMPILE__
+#  ifndef __gfx908__
+#    error "Passed architecture gfx908, but got something else."
+#  endif
+#endif
+
+int main()
+{
+}
diff --git a/Tests/HIP/CMakeLists.txt b/Tests/HIP/CMakeLists.txt
new file mode 100644
index 0000000..9499be8
--- /dev/null
+++ b/Tests/HIP/CMakeLists.txt
@@ -0,0 +1,15 @@
+macro (add_hip_test_macro name)
+  add_test_macro("${name}" ${ARGN})
+  set_property(TEST "${name}" APPEND
+    PROPERTY LABELS "HIP")
+endmacro ()
+
+add_hip_test_macro(HIP.ArchitectureOff HIPOnlyArchitectureOff)
+add_hip_test_macro(HIP.CompileFlags HIPOnlyCompileFlags)
+add_hip_test_macro(HIP.EnableStandard HIPEnableStandard)
+add_hip_test_macro(HIP.InferHipLang1 HIPInferHipLang1)
+add_hip_test_macro(HIP.InferHipLang2 HIPInferHipLang2)
+add_hip_test_macro(HIP.MathFunctions HIPOnlyMathFunctions)
+add_hip_test_macro(HIP.MixedLanguage HIPMixedLanguage)
+add_hip_test_macro(HIP.TryCompile HIPOnlyTryCompile)
+add_hip_test_macro(HIP.WithDefs HIPOnlyWithDefs)
diff --git a/Tests/HIP/CompileFlags/CMakeLists.txt b/Tests/HIP/CompileFlags/CMakeLists.txt
new file mode 100644
index 0000000..c808313
--- /dev/null
+++ b/Tests/HIP/CompileFlags/CMakeLists.txt
@@ -0,0 +1,8 @@
+cmake_minimum_required(VERSION 3.18)
+project(CompileFlags HIP)
+
+add_executable(HIPOnlyCompileFlags main.hip)
+
+set_property(TARGET HIPOnlyCompileFlags PROPERTY HIP_ARCHITECTURES gfx803)
+
+target_compile_options(HIPOnlyCompileFlags PRIVATE -DALWAYS_DEFINE)
diff --git a/Tests/HIP/CompileFlags/main.hip b/Tests/HIP/CompileFlags/main.hip
new file mode 100644
index 0000000..3259f1d
--- /dev/null
+++ b/Tests/HIP/CompileFlags/main.hip
@@ -0,0 +1,13 @@
+#ifdef __HIP_DEVICE_COMPILE__
+#  ifndef __gfx803__
+#    error "Passed architecture gfx803, but got something else."
+#  endif
+#endif
+
+#ifndef ALWAYS_DEFINE
+#  error "ALWAYS_DEFINE not defined!"
+#endif
+
+int main()
+{
+}
diff --git a/Tests/HIP/EnableStandard/CMakeLists.txt b/Tests/HIP/EnableStandard/CMakeLists.txt
new file mode 100644
index 0000000..6701724
--- /dev/null
+++ b/Tests/HIP/EnableStandard/CMakeLists.txt
@@ -0,0 +1,20 @@
+cmake_minimum_required(VERSION 3.18)
+project (EnableStandard HIP)
+
+set(CMAKE_CXX_COMPILER ${CMAKE_HIP_COMPILER})
+enable_language(CXX)
+
+#Goal for this example:
+#build hip sources that require C++11 to be enabled.
+
+add_library(HIPStatic11 STATIC static.cxx)
+set_source_files_properties(static.cxx PROPERTIES LANGUAGE HIP)
+
+add_library(HIPDynamic11 SHARED shared.hip)
+
+add_executable(HIPEnableStandard main.hip)
+target_link_libraries(HIPEnableStandard PRIVATE HIPStatic11 HIPDynamic11)
+
+target_compile_features(HIPDynamic11 PRIVATE cxx_std_11)
+set_target_properties(HIPStatic11 PROPERTIES HIP_STANDARD 11)
+set_target_properties(HIPStatic11 PROPERTIES HIP_STANDARD_REQUIRED TRUE)
diff --git a/Tests/HIP/EnableStandard/main.hip b/Tests/HIP/EnableStandard/main.hip
new file mode 100644
index 0000000..7dc32e8
--- /dev/null
+++ b/Tests/HIP/EnableStandard/main.hip
@@ -0,0 +1,23 @@
+
+#include <iostream>
+
+#ifdef _WIN32
+#  define IMPORT __declspec(dllimport)
+#else
+#  define IMPORT
+#endif
+
+int static_hip11_func(int);
+IMPORT int shared_hip11_func(int);
+
+void test_functions()
+{
+  static_hip11_func(int(42));
+  shared_hip11_func(int(42));
+}
+
+int main(int argc, char** argv)
+{
+  test_functions();
+  return 0;
+}
diff --git a/Tests/HIP/EnableStandard/shared.hip b/Tests/HIP/EnableStandard/shared.hip
new file mode 100644
index 0000000..2042258
--- /dev/null
+++ b/Tests/HIP/EnableStandard/shared.hip
@@ -0,0 +1,15 @@
+
+#include <type_traits>
+
+#ifdef _WIN32
+#  define EXPORT __declspec(dllexport)
+#else
+#  define EXPORT
+#endif
+
+using tt = std::true_type;
+using ft = std::false_type;
+EXPORT int __host__ shared_hip11_func(int x)
+{
+  return x * x + std::integral_constant<int, 17>::value;
+}
diff --git a/Tests/HIP/EnableStandard/static.cxx b/Tests/HIP/EnableStandard/static.cxx
new file mode 100644
index 0000000..2955894
--- /dev/null
+++ b/Tests/HIP/EnableStandard/static.cxx
@@ -0,0 +1,9 @@
+
+#include <type_traits>
+
+using tt = std::true_type;
+using ft = std::false_type;
+int __host__ static_hip11_func(int x)
+{
+  return x * x + std::integral_constant<int, 17>::value;
+}
diff --git a/Tests/HIP/InferHipLang1/CMakeLists.txt b/Tests/HIP/InferHipLang1/CMakeLists.txt
new file mode 100644
index 0000000..63d77fd
--- /dev/null
+++ b/Tests/HIP/InferHipLang1/CMakeLists.txt
@@ -0,0 +1,12 @@
+cmake_minimum_required(VERSION 3.18)
+project(InferHipLang C CXX HIP)
+
+#Goal for this example:
+#make sure that we understand that HIP is the correct link language
+add_library(InterfaceWithHIP INTERFACE)
+target_sources(InterfaceWithHIP INTERFACE interface.hip main.cxx)
+target_compile_features(InterfaceWithHIP INTERFACE hip_std_14)
+target_compile_features(InterfaceWithHIP INTERFACE cxx_std_11)
+
+add_executable(HIPInferHipLang1 )
+target_link_libraries(HIPInferHipLang1 PRIVATE InterfaceWithHIP)
diff --git a/Tests/HIP/InferHipLang1/interface.hip b/Tests/HIP/InferHipLang1/interface.hip
new file mode 100644
index 0000000..6ac8641
--- /dev/null
+++ b/Tests/HIP/InferHipLang1/interface.hip
@@ -0,0 +1,19 @@
+#include <system_error>
+#include <type_traits>
+#include <hip/hip_runtime_api.h>
+
+static __global__ void fake_hip_kernel()
+{
+}
+
+
+int __host__ interface_hip_func(int x)
+{
+
+  fake_hip_kernel<<<1, 1>>>();
+  bool valid = (hipSuccess == hipGetLastError());
+  if (!valid) {
+    throw std::system_error(ENODEV, std::generic_category(), "no hip device");
+  }
+  return x * x + std::integral_constant<int, 17>::value;
+}
diff --git a/Tests/HIP/InferHipLang1/main.cxx b/Tests/HIP/InferHipLang1/main.cxx
new file mode 100644
index 0000000..987b6c6
--- /dev/null
+++ b/Tests/HIP/InferHipLang1/main.cxx
@@ -0,0 +1,19 @@
+
+#include <iostream>
+
+#ifdef __HIP_PLATFORM_HCC__
+#  error "__HIP_PLATFORM_HCC__ propagated to C++ compilation!"
+#endif
+
+#ifdef __HIP_ROCclr__
+#  error "__HIP_ROCclr__ propagated to C++ compilation!"
+#endif
+
+int interface_hip_func(int);
+
+int main(int argc, char** argv)
+{
+  interface_hip_func(int(42));
+
+  return 0;
+}
diff --git a/Tests/HIP/InferHipLang2/CMakeLists.txt b/Tests/HIP/InferHipLang2/CMakeLists.txt
new file mode 100644
index 0000000..0e69de3
--- /dev/null
+++ b/Tests/HIP/InferHipLang2/CMakeLists.txt
@@ -0,0 +1,12 @@
+cmake_minimum_required(VERSION 3.18)
+project(InferHipLang C CXX HIP)
+
+#Goal for this example:
+#make sure that we understand that HIP is the correct link language
+add_library(InterfaceWithHIP OBJECT)
+target_sources(InterfaceWithHIP PRIVATE interface.hip main.cxx)
+target_compile_features(InterfaceWithHIP INTERFACE hip_std_14)
+target_compile_features(InterfaceWithHIP INTERFACE cxx_std_11)
+
+add_executable(HIPInferHipLang2 )
+target_link_libraries(HIPInferHipLang2 PRIVATE InterfaceWithHIP)
diff --git a/Tests/HIP/InferHipLang2/interface.hip b/Tests/HIP/InferHipLang2/interface.hip
new file mode 100644
index 0000000..6ac8641
--- /dev/null
+++ b/Tests/HIP/InferHipLang2/interface.hip
@@ -0,0 +1,19 @@
+#include <system_error>
+#include <type_traits>
+#include <hip/hip_runtime_api.h>
+
+static __global__ void fake_hip_kernel()
+{
+}
+
+
+int __host__ interface_hip_func(int x)
+{
+
+  fake_hip_kernel<<<1, 1>>>();
+  bool valid = (hipSuccess == hipGetLastError());
+  if (!valid) {
+    throw std::system_error(ENODEV, std::generic_category(), "no hip device");
+  }
+  return x * x + std::integral_constant<int, 17>::value;
+}
diff --git a/Tests/HIP/InferHipLang2/main.cxx b/Tests/HIP/InferHipLang2/main.cxx
new file mode 100644
index 0000000..987b6c6
--- /dev/null
+++ b/Tests/HIP/InferHipLang2/main.cxx
@@ -0,0 +1,19 @@
+
+#include <iostream>
+
+#ifdef __HIP_PLATFORM_HCC__
+#  error "__HIP_PLATFORM_HCC__ propagated to C++ compilation!"
+#endif
+
+#ifdef __HIP_ROCclr__
+#  error "__HIP_ROCclr__ propagated to C++ compilation!"
+#endif
+
+int interface_hip_func(int);
+
+int main(int argc, char** argv)
+{
+  interface_hip_func(int(42));
+
+  return 0;
+}
diff --git a/Tests/HIP/MathFunctions/CMakeLists.txt b/Tests/HIP/MathFunctions/CMakeLists.txt
new file mode 100644
index 0000000..81e3ddb
--- /dev/null
+++ b/Tests/HIP/MathFunctions/CMakeLists.txt
@@ -0,0 +1,18 @@
+cmake_minimum_required(VERSION 3.18)
+project(MathFunctions HIP)
+
+# This test covers these major HIP language/runtime requirements:
+#
+# 1. This makes sure CMake properly specifies the internal clang header dirs
+#    that hold headers needed for overloads of device side functions
+#
+# 2. This makes sure that all HIP include directories are properly marked as
+#    system includes so we don't get the following warnings:
+#      replacement function 'operator delete' cannot be declared 'inline'#
+#
+# 3. This makes sure CMake properly links to all the built-in libraries
+#    that hip needs that inject support for __half support
+#
+add_executable(HIPOnlyMathFunctions main.hip)
+target_compile_options(HIPOnlyMathFunctions PRIVATE -Werror)
+target_compile_features(HIPOnlyMathFunctions PRIVATE hip_std_14)
diff --git a/Tests/HIP/MathFunctions/main.hip b/Tests/HIP/MathFunctions/main.hip
new file mode 100644
index 0000000..8a6e77f
--- /dev/null
+++ b/Tests/HIP/MathFunctions/main.hip
@@ -0,0 +1,40 @@
+
+#include <stdexcept>
+#include <cmath>
+#include <math.h>
+#include <memory>
+
+#include <hip/hip_runtime.h>
+#include <hip/hip_fp16.h>
+
+namespace {
+template<class T, class F>
+__global__ void global_entry_point(F f, T *out) {
+  *out = f();
+}
+
+template <class T, class F>
+bool verify(F f, T expected)
+{
+  std::unique_ptr<T> cpu_T(new T);
+  T* gpu_T = nullptr;
+  hipMalloc((void**)&gpu_T, sizeof(T));
+  hipLaunchKernelGGL(global_entry_point, 1, 1, 0, 0, f, gpu_T);
+  hipMemcpy(cpu_T.get(), gpu_T, sizeof(T), hipMemcpyDeviceToHost);
+  hipFree(gpu_T);
+  return (*cpu_T == expected);
+}
+}
+
+int main(int argc, char** argv)
+{
+  bool valid = verify([]__device__(){ return std::round(1.4f); }, 1.0f);
+  valid &= verify([]__device__(){ return max<_Float16>(1.0f, 2.0f); }, 2.0f);
+  valid &= verify([]__device__(){ return min<_Float16>(1.0f, 2.0f); }, 1.0f);
+
+  if (valid) {
+    return 0;
+  } else {
+    return 1;
+  }
+}
diff --git a/Tests/HIP/MixedLanguage/CMakeLists.txt b/Tests/HIP/MixedLanguage/CMakeLists.txt
new file mode 100644
index 0000000..4f6dd3b
--- /dev/null
+++ b/Tests/HIP/MixedLanguage/CMakeLists.txt
@@ -0,0 +1,19 @@
+cmake_minimum_required(VERSION 3.18)
+project (MixedLanguage C CXX HIP)
+
+set(CMAKE_HIP_STANDARD 14)
+set(CMAKE_CXX_STANDARD 14)
+
+#Goal for this example:
+#make sure that we can build multiple languages into targets
+#and have the link language always be HIP
+add_library(MixedSharedLib SHARED shared.c)
+add_library(MixedObjectLib OBJECT shared.cxx shared.hip)
+set_target_properties(MixedObjectLib PROPERTIES POSITION_INDEPENDENT_CODE ON)
+target_link_libraries(MixedSharedLib PRIVATE MixedObjectLib)
+
+add_library(MixedStaticLib STATIC static.c static.cxx static.hip)
+set_target_properties(MixedStaticLib PROPERTIES POSITION_INDEPENDENT_CODE ON)
+
+add_executable(HIPMixedLanguage main.cxx)
+target_link_libraries(HIPMixedLanguage PRIVATE MixedStaticLib MixedSharedLib)
diff --git a/Tests/HIP/MixedLanguage/main.cxx b/Tests/HIP/MixedLanguage/main.cxx
new file mode 100644
index 0000000..003d18a
--- /dev/null
+++ b/Tests/HIP/MixedLanguage/main.cxx
@@ -0,0 +1,40 @@
+
+#include <iostream>
+
+#ifdef __HIP_PLATFORM_HCC__
+#  error "__HIP_PLATFORM_HCC__ propagated to C++ compilation!"
+#endif
+
+#ifdef __HIP_ROCclr__
+#  error "__HIP_ROCclr__ propagated to C++ compilation!"
+#endif
+
+#ifdef _WIN32
+#  define IMPORT __declspec(dllimport)
+#else
+#  define IMPORT
+#endif
+
+extern "C" {
+IMPORT int shared_c_func(int);
+int static_c_func(int);
+}
+
+IMPORT int shared_cxx_func(int);
+IMPORT int shared_hip_func(int);
+
+int static_cxx_func(int);
+int static_hip_func(int);
+
+int main(int argc, char** argv)
+{
+  static_c_func(int(42));
+  static_cxx_func(int(42));
+  static_hip_func(int(42));
+
+  shared_c_func(int(42));
+  shared_cxx_func(int(42));
+  shared_hip_func(int(42));
+
+  return 0;
+}
diff --git a/Tests/HIP/MixedLanguage/shared.c b/Tests/HIP/MixedLanguage/shared.c
new file mode 100644
index 0000000..347fba9
--- /dev/null
+++ b/Tests/HIP/MixedLanguage/shared.c
@@ -0,0 +1,12 @@
+
+
+#ifdef _WIN32
+#  define EXPORT __declspec(dllexport)
+#else
+#  define EXPORT
+#endif
+
+EXPORT int shared_c_func(int x)
+{
+  return -x;
+}
diff --git a/Tests/HIP/MixedLanguage/shared.cxx b/Tests/HIP/MixedLanguage/shared.cxx
new file mode 100644
index 0000000..8e6c1d3
--- /dev/null
+++ b/Tests/HIP/MixedLanguage/shared.cxx
@@ -0,0 +1,21 @@
+
+#include <type_traits>
+
+#ifdef __HIP_PLATFORM_HCC__
+#  error "__HIP_PLATFORM_HCC__ propagated to C++ compilation!"
+#endif
+
+#ifdef __HIP_ROCclr__
+#  error "__HIP_ROCclr__ propagated to C++ compilation!"
+#endif
+
+#ifdef _WIN32
+#  define EXPORT __declspec(dllexport)
+#else
+#  define EXPORT
+#endif
+
+EXPORT int shared_cxx_func(int x)
+{
+  return x * x + std::integral_constant<int, 14>::value;
+}
diff --git a/Tests/HIP/MixedLanguage/shared.hip b/Tests/HIP/MixedLanguage/shared.hip
new file mode 100644
index 0000000..e6fea9f
--- /dev/null
+++ b/Tests/HIP/MixedLanguage/shared.hip
@@ -0,0 +1,26 @@
+#include <system_error>
+#include <type_traits>
+#include <hip/hip_runtime_api.h>
+
+#ifdef _WIN32
+#  define EXPORT __declspec(dllexport)
+#else
+#  define EXPORT
+#endif
+
+
+static __global__ void fake_hip_kernel()
+{
+}
+
+
+int __host__ shared_hip_func(int x)
+{
+
+  fake_hip_kernel<<<1, 1>>>();
+  bool valid = (hipSuccess == hipGetLastError());
+  if (!valid) {
+    throw std::system_error(ENODEV, std::generic_category(), "no hip device");
+  }
+  return x * x + std::integral_constant<int, 17>::value;
+}
diff --git a/Tests/HIP/MixedLanguage/static.c b/Tests/HIP/MixedLanguage/static.c
new file mode 100644
index 0000000..06c33b4
--- /dev/null
+++ b/Tests/HIP/MixedLanguage/static.c
@@ -0,0 +1,6 @@
+
+
+int static_c_func(int x)
+{
+  return -x;
+}
diff --git a/Tests/HIP/MixedLanguage/static.cxx b/Tests/HIP/MixedLanguage/static.cxx
new file mode 100644
index 0000000..2c14fb1
--- /dev/null
+++ b/Tests/HIP/MixedLanguage/static.cxx
@@ -0,0 +1,7 @@
+
+#include <type_traits>
+
+int static_cxx_func(int x)
+{
+  return x * x + std::integral_constant<int, 14>::value;
+}
diff --git a/Tests/HIP/MixedLanguage/static.hip b/Tests/HIP/MixedLanguage/static.hip
new file mode 100644
index 0000000..359b9fa
--- /dev/null
+++ b/Tests/HIP/MixedLanguage/static.hip
@@ -0,0 +1,21 @@
+
+#include <type_traits>
+#include <system_error>
+#include <hip/hip_runtime_api.h>
+
+
+static __global__ void fake_hip_kernel()
+{
+}
+
+
+int __host__ static_hip_func(int x)
+{
+
+  fake_hip_kernel<<<1, 1>>>();
+  bool valid = (hipSuccess == hipGetLastError());
+  if (!valid) {
+    throw std::system_error(ENODEV, std::generic_category(), "no hip device");
+  }
+  return x * x + std::integral_constant<int, 17>::value;
+}
diff --git a/Tests/HIP/TryCompile/CMakeLists.txt b/Tests/HIP/TryCompile/CMakeLists.txt
new file mode 100644
index 0000000..92a834c
--- /dev/null
+++ b/Tests/HIP/TryCompile/CMakeLists.txt
@@ -0,0 +1,15 @@
+cmake_minimum_required(VERSION 3.18)
+project (TryCompile HIP)
+
+#Goal for this example:
+# Verify try_compile with HIP language works
+set(CMAKE_HIP_STANDARD 14)
+set(CMAKE_HIP_ARCHITECTURES gfx803 gfx900)
+
+set(CMAKE_TRY_COMPILE_TARGET_TYPE STATIC_LIBRARY)
+try_compile(result "${CMAKE_CURRENT_BINARY_DIR}"
+        SOURCES "${CMAKE_CURRENT_SOURCE_DIR}/device_function.hip"
+        COPY_FILE "${CMAKE_CURRENT_BINARY_DIR}/device_function.o")
+
+add_executable(HIPOnlyTryCompile main.hip)
+target_link_libraries(HIPOnlyTryCompile "${CMAKE_CURRENT_BINARY_DIR}/device_function.o")
diff --git a/Tests/HIP/TryCompile/device_function.hip b/Tests/HIP/TryCompile/device_function.hip
new file mode 100644
index 0000000..7c1205e
--- /dev/null
+++ b/Tests/HIP/TryCompile/device_function.hip
@@ -0,0 +1,17 @@
+#include <system_error>
+#include <hip/hip_runtime_api.h>
+
+static __global__ void fake_hip_kernel()
+{
+}
+
+int __host__ try_compile_hip_func(int x)
+{
+
+  fake_hip_kernel<<<1, 1>>>();
+  bool valid = (hipSuccess == hipGetLastError());
+  if (!valid) {
+    throw std::system_error(ENODEV, std::generic_category(), "no hip device");
+  }
+  return x * x;
+}
diff --git a/Tests/HIP/TryCompile/main.hip b/Tests/HIP/TryCompile/main.hip
new file mode 100644
index 0000000..091dca3
--- /dev/null
+++ b/Tests/HIP/TryCompile/main.hip
@@ -0,0 +1,8 @@
+int __host__ try_compile_hip_func(int x);
+
+int main(int argc, char** argv)
+{
+  try_compile_hip_func(int(42));
+
+  return 0;
+}
diff --git a/Tests/HIP/WithDefs/CMakeLists.txt b/Tests/HIP/WithDefs/CMakeLists.txt
new file mode 100644
index 0000000..e2db182
--- /dev/null
+++ b/Tests/HIP/WithDefs/CMakeLists.txt
@@ -0,0 +1,37 @@
+
+cmake_minimum_required(VERSION 3.18)
+project (WithDefs HIP)
+
+set(CMAKE_HIP_ARCHITECTURES OFF)
+set(release_compile_defs DEFREL)
+
+#Goal for this example:
+#build a executable that needs to be passed a complex define through add_definitions
+#this verifies we can pass C++ style attributes to hipcc
+add_definitions("-DPACKED_DEFINE=[[gnu::packed]]")
+
+add_executable(HIPOnlyWithDefs main.hip.cpp)
+set_source_files_properties(main.hip.cpp PROPERTIES LANGUAGE HIP)
+
+target_compile_features(HIPOnlyWithDefs PRIVATE hip_std_17)
+
+target_compile_options(HIPOnlyWithDefs
+  PRIVATE
+    --offload-arch=gfx900
+    -DFLAG_COMPILE_LANG_$<COMPILE_LANGUAGE>
+    $<$<HIP_COMPILER_ID:ROCMClang>:-DFLAG_LANG_IS_HIP=$<COMPILE_LANGUAGE:HIP>> # Host-only defines are possible only on NVCC.
+  )
+
+target_compile_definitions(HIPOnlyWithDefs
+  PRIVATE
+    $<$<CONFIG:RELEASE>:$<BUILD_INTERFACE:${release_compile_defs}>>
+    -DDEF_COMPILE_LANG_$<COMPILE_LANGUAGE>
+    -DDEF_LANG_IS_HIP=$<COMPILE_LANGUAGE:HIP>
+    -DDEF_HIP_COMPILER=$<HIP_COMPILER_ID>
+    -DDEF_HIP_COMPILER_VERSION=$<HIP_COMPILER_VERSION>
+  )
+
+target_include_directories(HIPOnlyWithDefs
+  PRIVATE
+    $<$<COMPILE_LANGUAGE:HIP>:${CMAKE_CURRENT_SOURCE_DIR}/inc_hip>
+)
diff --git a/Tests/HIP/WithDefs/inc_hip/inc_hip.h b/Tests/HIP/WithDefs/inc_hip/inc_hip.h
new file mode 100644
index 0000000..fe2698a
--- /dev/null
+++ b/Tests/HIP/WithDefs/inc_hip/inc_hip.h
@@ -0,0 +1 @@
+#define INC_HIP
diff --git a/Tests/HIP/WithDefs/main.hip.cpp b/Tests/HIP/WithDefs/main.hip.cpp
new file mode 100644
index 0000000..a8f2d18
--- /dev/null
+++ b/Tests/HIP/WithDefs/main.hip.cpp
@@ -0,0 +1,89 @@
+#include <iostream>
+
+#include <hip/hip_runtime_api.h>
+#include <inc_hip.h>
+#ifndef INC_HIP
+#  error "INC_HIP not defined!"
+#endif
+
+#ifndef PACKED_DEFINE
+#  error "PACKED_DEFINE not defined!"
+#endif
+
+#ifndef FLAG_COMPILE_LANG_HIP
+#  error "FLAG_COMPILE_LANG_HIP not defined!"
+#endif
+
+#ifndef FLAG_LANG_IS_HIP
+#  error "FLAG_LANG_IS_HIP not defined!"
+#endif
+
+#if !FLAG_LANG_IS_HIP
+#  error "Expected FLAG_LANG_IS_HIP"
+#endif
+
+#ifndef DEF_COMPILE_LANG_HIP
+#  error "DEF_COMPILE_LANG_HIP not defined!"
+#endif
+
+#ifndef DEF_LANG_IS_HIP
+#  error "DEF_LANG_IS_HIP not defined!"
+#endif
+
+#if !DEF_LANG_IS_HIP
+#  error "Expected DEF_LANG_IS_HIP"
+#endif
+
+#ifndef DEF_HIP_COMPILER
+#  error "DEF_HIP_COMPILER not defined!"
+#endif
+
+#ifndef DEF_HIP_COMPILER_VERSION
+#  error "DEF_HIP_COMPILER_VERSION not defined!"
+#endif
+
+static __global__ void DetermineIfValidHIPDevice()
+{
+}
+
+#ifdef _MSC_VER
+#  pragma pack(push, 1)
+#  undef PACKED_DEFINE
+#  define PACKED_DEFINE
+#endif
+struct PACKED_DEFINE result_type
+{
+  bool valid;
+  int value;
+#if defined(NDEBUG) && !defined(DEFREL)
+#  error missing DEFREL flag
+#endif
+};
+#ifdef _MSC_VER
+#  pragma pack(pop)
+#endif
+
+result_type can_launch_kernel()
+{
+  result_type r;
+  DetermineIfValidHIPDevice<<<1, 1>>>();
+  r.valid = (hipSuccess == hipGetLastError());
+  if (r.valid) {
+    r.value = 1;
+  } else {
+    r.value = -1;
+  }
+  return r;
+}
+
+int main(int argc, char** argv)
+{
+  hipError_t err;
+  int nDevices = 0;
+  err = hipGetDeviceCount(&nDevices);
+  if (err != hipSuccess) {
+    std::cerr << hipGetErrorString(err) << std::endl;
+    return 1;
+  }
+  return 0;
+}
diff --git a/Tests/RunCMake/CMakeLists.txt b/Tests/RunCMake/CMakeLists.txt
index 670abbc..4e28549 100644
--- a/Tests/RunCMake/CMakeLists.txt
+++ b/Tests/RunCMake/CMakeLists.txt
@@ -452,6 +452,7 @@
       CMAKE_CXX_STANDARD_DEFAULT
       CMake_TEST_CUDA
       CMake_TEST_ISPC
+      CMake_TEST_HIP
       CMake_TEST_FILESYSTEM_1S
       CMAKE_OBJC_STANDARD_DEFAULT
       CMAKE_OBJCXX_STANDARD_DEFAULT
@@ -603,12 +604,15 @@
 add_RunCMake_test(target_sources)
 add_RunCMake_test(CheckCompilerFlag   -DCMake_TEST_CUDA=${CMake_TEST_CUDA}
                                       -DCMake_TEST_ISPC=${CMake_TEST_ISPC}
-                                      -DCMAKE_Fortran_COMPILER_ID=${CMAKE_Fortran_COMPILER_ID})
+                                      -DCMAKE_Fortran_COMPILER_ID=${CMAKE_Fortran_COMPILER_ID}
+                                      -DCMake_TEST_HIP=${CMake_TEST_HIP})
 add_RunCMake_test(CheckSourceCompiles -DCMake_TEST_CUDA=${CMake_TEST_CUDA}
                                       -DCMake_TEST_ISPC=${CMake_TEST_ISPC}
-                                      -DCMAKE_Fortran_COMPILER_ID=${CMAKE_Fortran_COMPILER_ID})
+                                      -DCMAKE_Fortran_COMPILER_ID=${CMAKE_Fortran_COMPILER_ID}
+                                      -DCMake_TEST_HIP=${CMake_TEST_HIP})
 add_RunCMake_test(CheckSourceRuns     -DCMake_TEST_CUDA=${CMake_TEST_CUDA}
-                                      -DCMAKE_Fortran_COMPILER_ID=${CMAKE_Fortran_COMPILER_ID})
+                                      -DCMAKE_Fortran_COMPILER_ID=${CMAKE_Fortran_COMPILER_ID}
+                                      -DCMake_TEST_HIP=${CMake_TEST_HIP})
 set_property(TEST RunCMake.CheckCompilerFlag
                   RunCMake.CheckSourceCompiles
                   RunCMake.CheckSourceRuns
@@ -616,14 +620,20 @@
 set_property(TEST RunCMake.CheckSourceCompiles
                   RunCMake.CheckCompilerFlag
     APPEND PROPERTY LABELS "ISPC")
+set_property(TEST RunCMake.CheckCompilerFlag
+                  RunCMake.CheckSourceCompiles
+                  RunCMake.CheckSourceRuns
+    APPEND PROPERTY LABELS "HIP")
 add_RunCMake_test(CheckModules)
 add_RunCMake_test(CheckIPOSupported)
 if (CMAKE_SYSTEM_NAME MATCHES "(Linux|Darwin)"
     AND (CMAKE_C_COMPILER_ID MATCHES "Clang|GNU" OR CMAKE_Fortran_COMPILER_ID MATCHES "GNU"))
   add_RunCMake_test(CheckLinkerFlag -DCMAKE_C_COMPILER_ID=${CMAKE_C_COMPILER_ID}
                                     -DCMAKE_Fortran_COMPILER_ID=${CMAKE_Fortran_COMPILER_ID}
-                                    -DCMake_TEST_CUDA=${CMake_TEST_CUDA})
+                                    -DCMake_TEST_CUDA=${CMake_TEST_CUDA}
+                                    -DCMake_TEST_HIP=${CMake_TEST_HIP})
   set_property(TEST RunCMake.CheckLinkerFlag APPEND PROPERTY LABELS "CUDA")
+  set_property(TEST RunCMake.CheckLinkerFlag APPEND PROPERTY LABELS "HIP")
 endif()
 
 
@@ -727,6 +737,9 @@
   if(DEFINED CMake_TEST_CUDA)
     list(APPEND CompilerLauncher_ARGS -DCMake_TEST_CUDA=${CMake_TEST_CUDA})
   endif()
+  if(DEFINED CMake_TEST_HIP)
+    list(APPEND CompilerLauncher_ARGS -DCMake_TEST_HIP=${CMake_TEST_HIP})
+  endif()
   if(DEFINED CMake_TEST_ISPC)
     list(APPEND CompilerLauncher_ARGS -DCMake_TEST_ISPC=${CMake_TEST_ISPC})
   endif()
@@ -739,7 +752,7 @@
   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)
@@ -839,9 +852,19 @@
   -DCMAKE_C_COMPILER_VERSION=${CMAKE_C_COMPILER_VERSION})
 
 add_RunCMake_test("UnityBuild")
-add_RunCMake_test(CMakePresets -DPYTHON_EXECUTABLE=${PYTHON_EXECUTABLE} -DCMake_TEST_JSON_SCHEMA=${CMake_TEST_JSON_SCHEMA})
-add_RunCMake_test(CMakePresetsBuild -DPYTHON_EXECUTABLE=${PYTHON_EXECUTABLE} -DCMake_TEST_JSON_SCHEMA=${CMake_TEST_JSON_SCHEMA})
-add_RunCMake_test(CMakePresetsTest -DPYTHON_EXECUTABLE=${PYTHON_EXECUTABLE} -DCMake_TEST_JSON_SCHEMA=${CMake_TEST_JSON_SCHEMA})
+add_RunCMake_test(CMakePresets
+  -DPYTHON_EXECUTABLE=${PYTHON_EXECUTABLE}
+  -DCMake_TEST_JSON_SCHEMA=${CMake_TEST_JSON_SCHEMA}
+  )
+add_RunCMake_test(CMakePresetsBuild
+  -DPYTHON_EXECUTABLE=${PYTHON_EXECUTABLE}
+  -DCMake_TEST_JSON_SCHEMA=${CMake_TEST_JSON_SCHEMA}
+  -DCMAKE_SYSTEM_NAME=${CMAKE_SYSTEM_NAME}
+  )
+add_RunCMake_test(CMakePresetsTest
+  -DPYTHON_EXECUTABLE=${PYTHON_EXECUTABLE}
+  -DCMake_TEST_JSON_SCHEMA=${CMake_TEST_JSON_SCHEMA}
+  )
 
 if(${CMAKE_GENERATOR} MATCHES "Make|Ninja")
   add_RunCMake_test(TransformDepfile)
diff --git a/Tests/RunCMake/CMakePresetsBuild/Good.json.in b/Tests/RunCMake/CMakePresetsBuild/Good.json.in
index fd43c4e..c7f318c 100644
--- a/Tests/RunCMake/CMakePresetsBuild/Good.json.in
+++ b/Tests/RunCMake/CMakePresetsBuild/Good.json.in
@@ -30,7 +30,7 @@
             "description": "",
             "inheritConfigureEnvironment": true,
             "environment": {},
-            "jobs": 0,
+            @Good_json_jobs@
             "targets": [],
             "configuration": "",
             "verbose": true,
@@ -73,6 +73,11 @@
             "vendor": {
                 "example.com": "value"
             }
+        },
+        {
+            "name": "singleTarget",
+            "inherits": "build-default",
+            "targets": "good"
         }
     ]
 }
diff --git a/Tests/RunCMake/CMakePresetsBuild/RunCMakeTest.cmake b/Tests/RunCMake/CMakePresetsBuild/RunCMakeTest.cmake
index afa22eb..b37c770 100644
--- a/Tests/RunCMake/CMakePresetsBuild/RunCMakeTest.cmake
+++ b/Tests/RunCMake/CMakePresetsBuild/RunCMakeTest.cmake
@@ -62,7 +62,15 @@
 
 set(CMakePresets_SCHEMA_EXPECTED_RESULT 0)
 
-run_cmake_build_presets(Good "default;other" "build-other;withEnvironment;noEnvironment;macros;vendorObject")
+if(RunCMake_GENERATOR MATCHES "NMake|Borland|Watcom")
+  set(Good_json_jobs [[]])
+elseif(RunCMake_GENERATOR MATCHES "Make" AND CMAKE_SYSTEM_NAME MATCHES "FreeBSD")
+  set(Good_json_jobs [["jobs": 1,]])
+else()
+  set(Good_json_jobs [["jobs": 0,]])
+endif()
+
+run_cmake_build_presets(Good "default;other" "build-other;withEnvironment;noEnvironment;macros;vendorObject;singleTarget")
 run_cmake_build_presets(InvalidConfigurePreset "default" "badConfigurePreset")
 run_cmake_build_presets(Condition "default" "enabled;disabled")
 
diff --git a/Tests/RunCMake/CPack/VerifyResult.cmake b/Tests/RunCMake/CPack/VerifyResult.cmake
index 294b9e8..76c1bea 100644
--- a/Tests/RunCMake/CPack/VerifyResult.cmake
+++ b/Tests/RunCMake/CPack/VerifyResult.cmake
@@ -87,7 +87,7 @@
           "which does not match:${msg_expected}\n"
           "${output_error_message}")
       endif()
-    elseif(foundFilescount_ EQUAL 0)
+    elseif(foundFilesCount_ EQUAL 0)
       message(FATAL_ERROR
         "Found no files for file ${file_no_}!\n"
         "Globbing expression:\n '${EXPECTED_FILE_${file_no_}}'\n"
diff --git a/Tests/RunCMake/CPack/tests/SOURCE_PACKAGE/VerifyResult.cmake b/Tests/RunCMake/CPack/tests/SOURCE_PACKAGE/VerifyResult.cmake
index 73d7481..42e8384 100644
--- a/Tests/RunCMake/CPack/tests/SOURCE_PACKAGE/VerifyResult.cmake
+++ b/Tests/RunCMake/CPack/tests/SOURCE_PACKAGE/VerifyResult.cmake
@@ -18,14 +18,20 @@
       OUTPUT_QUIET
   )
 
-set(output_error_message_
-    "\n${RPMBUILD_EXECUTABLE} error: '${error_}';\nresult: '${result_}';\n${output_error_message}")
+string(REPLACE "\n" "\n rpmbuild-err> " rpmbuild_err "\n${error_}")
+string(REPLACE "\n" "\n rpmbuild-res> " rpmbuild_res "\n${result_}")
+string(CONCAT output_error_message_
+  "${RPMBUILD_EXECUTABLE} error:${rpmbuild_err}\n"
+  "${RPMBUILD_EXECUTABLE} result:${rpmbuild_res}\n"
+  "${output_error_message}"
+  )
 
 # expected file content are test_prog and optional build-id links that are
 # generated by rpmbuild (introduced in rpm 4.13.0.1)
 set(EXPECTED_FILE_CONTENT_ "^/foo${whitespaces_}/foo/test_prog(${whitespaces_}.*\.build-id.*)*$")
+set(EXPECTED_FILE_ "${CMAKE_CURRENT_BINARY_DIR}/test_rpm/RPMS/*.rpm")
 
-file(GLOB_RECURSE FOUND_FILE_ RELATIVE "${CMAKE_CURRENT_BINARY_DIR}/test_rpm/RPMS" "${CMAKE_CURRENT_BINARY_DIR}/test_rpm/RPMS/*.rpm")
+file(GLOB_RECURSE FOUND_FILE_ RELATIVE "${CMAKE_CURRENT_BINARY_DIR}/test_rpm/RPMS" "${EXPECTED_FILE_}")
 list(APPEND foundFiles_ "${FOUND_FILE_}")
 list(LENGTH FOUND_FILE_ foundFilesCount_)
 
@@ -37,17 +43,27 @@
       expected_content_list "${PACKAGE_CONTENT}")
 
   if(NOT expected_content_list)
+    string(REPLACE "\n" "\n actual> " msg_actual "\n${PACKAGE_CONTENT}")
+    string(REPLACE "\n" "\n expect> " msg_expected "\n${EXPECTED_FILE_CONTENT_}")
+    string(REPLACE "\r" "\\r" msg_expected "${msg_expected}")
+    string(REPLACE "\t" "\\t" msg_expected "${msg_expected}")
     message(FATAL_ERROR
       "Unexpected file content!\n"
-      " Content: '${PACKAGE_CONTENT}'\n\n"
-      " Expected: '${EXPECTED_FILE_CONTENT_}'"
+      "The content was:${msg_actual}\n"
+      "which does not match:${msg_expected}\n"
       "${output_error_message_}")
   endif()
+elseif(foundFilesCount_ EQUAL 0)
+  message(FATAL_ERROR
+    "Found no files!\n"
+    "Globbing expression:\n '${EXPECTED_FILE_}'\n"
+    "${output_error_message_}")
 else()
   message(FATAL_ERROR
-    "Found more than one file!"
-    " Found files count '${foundFilesCount_}'."
-    " Files: '${FOUND_FILE_}'"
+    "Found more than one file!\n"
+    "Found files count '${foundFilesCount_}'.\n"
+    "Files:\n '${FOUND_FILE_}'\n"
+    "Globbing expression:\n '${EXPECTED_FILE_}'\n"
     "${output_error_message_}")
 endif()
 
@@ -62,6 +78,7 @@
 
 if(NOT foundFilesCount_ EQUAL allFoundFilesCount_)
   message(FATAL_ERROR
-      "Found more files than expected! Found files: '${allFoundFiles_}'"
-      "${output_error_message_}")
+    "Found more files than expected!\n"
+    "Found files:\n '${allFoundFiles_}'\n"
+    "${output_error_message_}")
 endif()
diff --git a/Tests/RunCMake/CheckCompilerFlag/CheckHIPCompilerFlag.cmake b/Tests/RunCMake/CheckCompilerFlag/CheckHIPCompilerFlag.cmake
new file mode 100644
index 0000000..339ce18
--- /dev/null
+++ b/Tests/RunCMake/CheckCompilerFlag/CheckHIPCompilerFlag.cmake
@@ -0,0 +1,13 @@
+
+enable_language (HIP)
+include(CheckCompilerFlag)
+
+check_compiler_flag(HIP "-_this_is_not_a_flag_" SHOULD_FAIL)
+if(SHOULD_FAIL)
+  message(SEND_ERROR "invalid HIP compile flag didn't fail.")
+endif()
+
+check_compiler_flag(HIP "-DFOO" SHOULD_WORK)
+if(NOT SHOULD_WORK)
+  message(SEND_ERROR "${CMAKE_HIP_COMPILER_ID} compiler flag '-DFOO' check failed")
+endif()
diff --git a/Tests/RunCMake/CheckCompilerFlag/RunCMakeTest.cmake b/Tests/RunCMake/CheckCompilerFlag/RunCMakeTest.cmake
index 7a4e2ce..7ef1860 100644
--- a/Tests/RunCMake/CheckCompilerFlag/RunCMakeTest.cmake
+++ b/Tests/RunCMake/CheckCompilerFlag/RunCMakeTest.cmake
@@ -22,3 +22,7 @@
 if(CMake_TEST_ISPC)
   run_cmake(CheckISPCCompilerFlag)
 endif()
+
+if(CMake_TEST_HIP)
+  run_cmake(CheckHIPCompilerFlag)
+endif()
diff --git a/Tests/RunCMake/CheckLinkerFlag/CheckHIPLinkerFlag.cmake b/Tests/RunCMake/CheckLinkerFlag/CheckHIPLinkerFlag.cmake
new file mode 100644
index 0000000..3bf3b30
--- /dev/null
+++ b/Tests/RunCMake/CheckLinkerFlag/CheckHIPLinkerFlag.cmake
@@ -0,0 +1,3 @@
+
+set (CHECK_LANGUAGE HIP)
+include ("${CMAKE_CURRENT_SOURCE_DIR}/CheckLinkerFlag.cmake")
diff --git a/Tests/RunCMake/CheckLinkerFlag/RunCMakeTest.cmake b/Tests/RunCMake/CheckLinkerFlag/RunCMakeTest.cmake
index 6ec9148..5e5bff6 100644
--- a/Tests/RunCMake/CheckLinkerFlag/RunCMakeTest.cmake
+++ b/Tests/RunCMake/CheckLinkerFlag/RunCMakeTest.cmake
@@ -16,3 +16,7 @@
 if (CMake_TEST_CUDA)
   run_cmake(CheckCUDALinkerFlag)
 endif()
+
+if (CMake_TEST_HIP)
+  run_cmake(CheckHIPLinkerFlag)
+endif()
diff --git a/Tests/RunCMake/CheckSourceCompiles/CheckHIPSourceCompiles.cmake b/Tests/RunCMake/CheckSourceCompiles/CheckHIPSourceCompiles.cmake
new file mode 100644
index 0000000..911a0d7
--- /dev/null
+++ b/Tests/RunCMake/CheckSourceCompiles/CheckHIPSourceCompiles.cmake
@@ -0,0 +1,27 @@
+
+enable_language (HIP)
+include(CheckSourceCompiles)
+
+check_source_compiles(HIP "I don't build" SHOULD_FAIL)
+if(SHOULD_FAIL)
+  message(SEND_ERROR "invalid HIP source didn't fail.")
+endif()
+
+check_source_compiles(HIP [=[
+  #include <vector>
+  __device__ int d_func() { }
+  int main() {
+    return 0;
+  }
+]=]
+ SHOULD_BUILD)
+if(NOT SHOULD_BUILD)
+  message(SEND_ERROR "Test fail for valid HIP source.")
+endif()
+
+check_source_compiles(HIP "void l(char const (&x)[2]){}; int main() { l(\"\\n\"); return 0;}"
+ SHOULD_BUILD_COMPLEX)
+
+if(NOT SHOULD_BUILD_COMPLEX)
+  message(SEND_ERROR "Test fail for valid HIP complex source.")
+endif()
diff --git a/Tests/RunCMake/CheckSourceCompiles/RunCMakeTest.cmake b/Tests/RunCMake/CheckSourceCompiles/RunCMakeTest.cmake
index 6e9088f..530f133 100644
--- a/Tests/RunCMake/CheckSourceCompiles/RunCMakeTest.cmake
+++ b/Tests/RunCMake/CheckSourceCompiles/RunCMakeTest.cmake
@@ -23,3 +23,7 @@
 if(CMake_TEST_ISPC)
   run_cmake(CheckISPCSourceCompiles)
 endif()
+
+if(CMake_TEST_HIP)
+  run_cmake(CheckHIPSourceCompiles)
+endif()
diff --git a/Tests/RunCMake/CheckSourceRuns/CheckHIPSourceRuns.cmake b/Tests/RunCMake/CheckSourceRuns/CheckHIPSourceRuns.cmake
new file mode 100644
index 0000000..d9fb8c2
--- /dev/null
+++ b/Tests/RunCMake/CheckSourceRuns/CheckHIPSourceRuns.cmake
@@ -0,0 +1,21 @@
+
+enable_language (HIP)
+include(CheckSourceRuns)
+
+check_source_runs(HIP "int main() {return 2;}" SHOULD_FAIL)
+if(SHOULD_FAIL)
+  message(SEND_ERROR "HIP check_source_runs succeeded, but should have failed.")
+endif()
+
+check_source_runs(HIP
+[=[
+  #include <vector>
+  __device__ __host__ void fake_function();
+  __host__ int main() {
+    return 0;
+  }
+]=]
+ SHOULD_RUN)
+if(NOT SHOULD_RUN)
+  message(SEND_ERROR "HIP check_source_runs failed for valid HIP executable.")
+endif()
diff --git a/Tests/RunCMake/CheckSourceRuns/RunCMakeTest.cmake b/Tests/RunCMake/CheckSourceRuns/RunCMakeTest.cmake
index c99ac8b..4784103 100644
--- a/Tests/RunCMake/CheckSourceRuns/RunCMakeTest.cmake
+++ b/Tests/RunCMake/CheckSourceRuns/RunCMakeTest.cmake
@@ -19,3 +19,7 @@
 if (CMake_TEST_CUDA)
   run_cmake(CheckCUDASourceRuns)
 endif()
+
+if (CMake_TEST_HIP)
+  run_cmake(CheckHIPSourceRuns)
+endif()
diff --git a/Tests/RunCMake/CompilerLauncher/HIP-common.cmake b/Tests/RunCMake/CompilerLauncher/HIP-common.cmake
new file mode 100644
index 0000000..53ece78
--- /dev/null
+++ b/Tests/RunCMake/CompilerLauncher/HIP-common.cmake
@@ -0,0 +1,5 @@
+enable_language(HIP)
+enable_language(CXX)
+set(CMAKE_VERBOSE_MAKEFILE TRUE)
+
+add_executable(main main.hip)
diff --git a/Tests/RunCMake/CompilerLauncher/HIP-env-Build-stdout.txt b/Tests/RunCMake/CompilerLauncher/HIP-env-Build-stdout.txt
new file mode 100644
index 0000000..3313e31
--- /dev/null
+++ b/Tests/RunCMake/CompilerLauncher/HIP-env-Build-stdout.txt
@@ -0,0 +1 @@
+.*-E env USED_LAUNCHER=1.*
diff --git a/Tests/RunCMake/CompilerLauncher/HIP-env-launch-Build-stdout.txt b/Tests/RunCMake/CompilerLauncher/HIP-env-launch-Build-stdout.txt
new file mode 100644
index 0000000..3313e31
--- /dev/null
+++ b/Tests/RunCMake/CompilerLauncher/HIP-env-launch-Build-stdout.txt
@@ -0,0 +1 @@
+.*-E env USED_LAUNCHER=1.*
diff --git a/Tests/RunCMake/CompilerLauncher/HIP-env.cmake b/Tests/RunCMake/CompilerLauncher/HIP-env.cmake
new file mode 100644
index 0000000..1bf56ce
--- /dev/null
+++ b/Tests/RunCMake/CompilerLauncher/HIP-env.cmake
@@ -0,0 +1 @@
+include(HIP-common.cmake)
diff --git a/Tests/RunCMake/CompilerLauncher/HIP-launch-Build-stdout.txt b/Tests/RunCMake/CompilerLauncher/HIP-launch-Build-stdout.txt
new file mode 100644
index 0000000..3313e31
--- /dev/null
+++ b/Tests/RunCMake/CompilerLauncher/HIP-launch-Build-stdout.txt
@@ -0,0 +1 @@
+.*-E env USED_LAUNCHER=1.*
diff --git a/Tests/RunCMake/CompilerLauncher/HIP-launch-env.cmake b/Tests/RunCMake/CompilerLauncher/HIP-launch-env.cmake
new file mode 100644
index 0000000..37985a5
--- /dev/null
+++ b/Tests/RunCMake/CompilerLauncher/HIP-launch-env.cmake
@@ -0,0 +1,3 @@
+set(CTEST_USE_LAUNCHERS 1)
+include(CTestUseLaunchers)
+include(HIP-env.cmake)
diff --git a/Tests/RunCMake/CompilerLauncher/HIP-launch.cmake b/Tests/RunCMake/CompilerLauncher/HIP-launch.cmake
new file mode 100644
index 0000000..78fd16b
--- /dev/null
+++ b/Tests/RunCMake/CompilerLauncher/HIP-launch.cmake
@@ -0,0 +1,3 @@
+set(CTEST_USE_LAUNCHERS 1)
+include(CTestUseLaunchers)
+include(HIP.cmake)
diff --git a/Tests/RunCMake/CompilerLauncher/HIP.cmake b/Tests/RunCMake/CompilerLauncher/HIP.cmake
new file mode 100644
index 0000000..9d2577a
--- /dev/null
+++ b/Tests/RunCMake/CompilerLauncher/HIP.cmake
@@ -0,0 +1,2 @@
+set(CMAKE_HIP_COMPILER_LAUNCHER "${CMAKE_COMMAND};-E;env;USED_LAUNCHER=1")
+include(HIP-common.cmake)
diff --git a/Tests/RunCMake/CompilerLauncher/RunCMakeTest.cmake b/Tests/RunCMake/CompilerLauncher/RunCMakeTest.cmake
index 787282a..84d0479 100644
--- a/Tests/RunCMake/CompilerLauncher/RunCMakeTest.cmake
+++ b/Tests/RunCMake/CompilerLauncher/RunCMakeTest.cmake
@@ -29,6 +29,9 @@
 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;
+}
diff --git a/Tests/RunCMake/ctest_test/RunCMakeTest.cmake b/Tests/RunCMake/ctest_test/RunCMakeTest.cmake
index b559e89..f07a12b 100644
--- a/Tests/RunCMake/ctest_test/RunCMakeTest.cmake
+++ b/Tests/RunCMake/ctest_test/RunCMakeTest.cmake
@@ -161,7 +161,23 @@
   NAME img_measurement
   COMMAND ${CMAKE_COMMAND} -E
   echo <DartMeasurementFile name="TestImage" type="image/png">]] ${IMAGE_DIR}/cmake-logo-16.png [[</DartMeasurementFile>)
+add_test(
+  NAME file_measurement
+  COMMAND ${CMAKE_COMMAND} -E
+  echo <DartMeasurementFile name="my_test_input_data" type="file">]] ${IMAGE_DIR}/cmake-logo-16.png [[</DartMeasurementFile>)
   ]])
   run_ctest(TestMeasurements)
 endfunction()
 run_measurements()
+
+# Verify that test output can override the Completion Status.
+function(run_completion_status)
+  set(CASE_CMAKELISTS_SUFFIX_CODE [[
+add_test(
+  NAME custom_details
+  COMMAND ${CMAKE_COMMAND} -E
+  echo test output\n<CTestDetails>CustomDetails</CTestDetails>\nmore output)
+  ]])
+  run_ctest(TestCompletionStatus)
+endfunction()
+run_completion_status()
diff --git a/Tests/RunCMake/ctest_test/TestCompletionStatus-check.cmake b/Tests/RunCMake/ctest_test/TestCompletionStatus-check.cmake
new file mode 100644
index 0000000..10de2ed
--- /dev/null
+++ b/Tests/RunCMake/ctest_test/TestCompletionStatus-check.cmake
@@ -0,0 +1,16 @@
+file(READ "${RunCMake_TEST_BINARY_DIR}/Testing/TAG" _tag)
+string(REGEX REPLACE "^([^\n]*)\n.*$" "\\1" _date "${_tag}")
+file(READ "${RunCMake_TEST_BINARY_DIR}/Testing/${_date}/Test.xml" _test_contents)
+
+# Check custom completion status.
+if(NOT _test_contents MATCHES [[<Value>CustomDetails</Value>]])
+  string(APPEND RunCMake_TEST_FAILED
+    "Could not find expected <Value>CustomDetails</Value> in Test.xml")
+endif()
+# Check test output.
+if(NOT _test_contents MATCHES "test output")
+  string(APPEND RunCMake_TEST_FAILED "Could not find expected string 'test output' in Test.xml")
+endif()
+if(NOT _test_contents MATCHES "more output")
+  string(APPEND RunCMake_TEST_FAILED "Could not find expected string 'more output' in Test.xml")
+endif()
diff --git a/Tests/RunCMake/ctest_test/TestMeasurements-check.cmake b/Tests/RunCMake/ctest_test/TestMeasurements-check.cmake
index 9ff9447..0095db0 100644
--- a/Tests/RunCMake/ctest_test/TestMeasurements-check.cmake
+++ b/Tests/RunCMake/ctest_test/TestMeasurements-check.cmake
@@ -15,3 +15,8 @@
   string(APPEND RunCMake_TEST_FAILED
     "Could not find expected <NamedMeasurement> tag for type='image/png' in Test.xml")
 endif()
+# Check file measurement.
+if(NOT _test_contents MATCHES [[NamedMeasurement name="my_test_input_data" encoding="base64" compression="tar/gzip" filename="cmake-logo-16.png" type="file"]])
+  string(APPEND RunCMake_TEST_FAILED
+    "Could not find expected <NamedMeasurement> tag for type='file' in Test.xml")
+endif()