diff options
| author | Brad King <brad.king@kitware.com> | 2021-06-09 11:53:16 (GMT) |
|---|---|---|
| committer | Kitware Robot <kwrobot@kitware.com> | 2021-06-09 11:53:32 (GMT) |
| commit | 7aad9e8685b0231abfb11d0f7d8c613cf5364fef (patch) | |
| tree | 2d2b547cbc6708936c98ae52c0f181d78fb0652d | |
| parent | b6faa5e2df22fecac877dc866294e8b5a36ab03a (diff) | |
| parent | 8514ee9b315b4ad02eed0e3cf11d8579f307fac0 (diff) | |
| download | CMake-7aad9e8685b0231abfb11d0f7d8c613cf5364fef.zip CMake-7aad9e8685b0231abfb11d0f7d8c613cf5364fef.tar.gz CMake-7aad9e8685b0231abfb11d0f7d8c613cf5364fef.tar.bz2 | |
Merge topic 'add_hip_language'
8514ee9b31 HIP: analyze output of `hipcc` to determine default GPU architecture
20d086f1a2 HIP: All HIP tests now run on CMake's current AMD hardware
2e86e50c2f HIP: Add HIP to all the Check* modules
947dbed0aa HIP: Automatically inject the `hip::device` runtime target
b50bfc8913 HIP: Add language to CMake
ff0d2858e1 HIP: Extract clang compiler details from hipcc
bd844387df ROCMClang: Add the ROCm toolkit derived clang compiler to CMake
590553f322 Compilers: protect use of __has_include
...
Acked-by: Kitware Robot <kwrobot@kitware.com>
Acked-by: Zack Galbreath <zack.galbreath@kitware.com>
Reviewed-by: Raul Tambre <raul@tambre.ee>
Acked-by: Axel Huebl <axel.huebl@plasma.ninja>
Merge-request: !6121
111 files changed, 2025 insertions, 66 deletions
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 @@ Enables support for the named language in CMake. This is the same as the :command:`project` command but does not create any of the extra variables that are created by the project command. Example languages are ``CXX``, ``C``, ``CUDA``, ``OBJC``, ``OBJCXX``, ``Fortran``, -``ISPC``, and ``ASM``. +``HIP``, ``ISPC``, and ``ASM``. .. versionadded:: 3.8 Added ``CUDA`` support. @@ -21,6 +21,9 @@ are ``CXX``, ``C``, ``CUDA``, ``OBJC``, ``OBJCXX``, ``Fortran``, .. versionadded:: 3.18 Added ``ISPC`` support. +.. versionadded:: 3.21 + Added ``HIP`` support. + If enabling ``ASM``, enable it last so that CMake can check whether compilers for other languages like ``C`` work for assembly too. 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 @@ The options are: Selects which programming languages are needed to build the project. Supported languages include ``C``, ``CXX`` (i.e. C++), ``CUDA``, - ``OBJC`` (i.e. Objective-C), ``OBJCXX``, ``Fortran``, ``ISPC``, and ``ASM``. + ``OBJC`` (i.e. Objective-C), ``OBJCXX``, ``Fortran``, ``HIP``, ``ISPC``, and ``ASM``. By default ``C`` and ``CXX`` are enabled if no language options are given. Specify language ``NONE``, or use the ``LANGUAGES`` keyword and list no languages, to skip enabling any languages. 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 @@ Environment Variables for Languages /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 @@ Variable Queries of the entries in ``compiler_ids``, otherwise ``0``. See also the :variable:`CMAKE_<LANG>_COMPILER_ID` variable. +.. genex:: $<HIP_COMPILER_ID:compiler_ids> + + where ``compiler_ids`` is a comma-separated list. + ``1`` if the CMake's compiler id of the HIP compiler matches any one + of the entries in ``compiler_ids``, otherwise ``0``. + See also the :variable:`CMAKE_<LANG>_COMPILER_ID` variable. + .. genex:: $<ISPC_COMPILER_ID:compiler_ids> where ``compiler_ids`` is a comma-separated list. @@ -218,6 +225,11 @@ Variable Queries ``1`` if the version of the Fortran compiler matches ``version``, otherwise ``0``. See also the :variable:`CMAKE_<LANG>_COMPILER_VERSION` variable. +.. genex:: $<HIP_COMPILER_VERSION:version> + + ``1`` if the version of the HIP compiler matches ``version``, otherwise ``0``. + See also the :variable:`CMAKE_<LANG>_COMPILER_VERSION` variable. + .. genex:: $<ISPC_COMPILER_VERSION:version> ``1`` if the version of the ISPC compiler matches ``version``, otherwise ``0``. @@ -648,6 +660,11 @@ Variable Queries The CMake's compiler id of the Fortran compiler used. See also the :variable:`CMAKE_<LANG>_COMPILER_ID` variable. +.. genex:: $<HIP_COMPILER_ID> + + The CMake's compiler id of the HIP compiler used. + See also the :variable:`CMAKE_<LANG>_COMPILER_ID` variable. + .. genex:: $<ISPC_COMPILER_ID> The CMake's compiler id of the ISPC compiler used. @@ -683,6 +700,11 @@ Variable Queries The version of the Fortran compiler used. See also the :variable:`CMAKE_<LANG>_COMPILER_VERSION` variable. +.. genex:: $<HIP_COMPILER_VERSION> + + The version of the HIP compiler used. + See also the :variable:`CMAKE_<LANG>_COMPILER_VERSION` variable. + .. genex:: $<ISPC_COMPILER_VERSION> The version of the ISPC compiler used. 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 @@ Properties on Targets /prop_tgt/GHS_NO_SOURCE_GROUP_FILE /prop_tgt/GNUtoMS /prop_tgt/HAS_CXX + /prop_tgt/HIP_ARCHITECTURES /prop_tgt/IMPLICIT_DEPENDS_INCLUDE_TRANSFORM /prop_tgt/IMPORTED /prop_tgt/IMPORTED_COMMON_LANGUAGE_RUNTIME 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 @@ Variables for Languages /variable/CMAKE_Fortran_MODDIR_DEFAULT /variable/CMAKE_Fortran_MODDIR_FLAG /variable/CMAKE_Fortran_MODOUT_FLAG + /variable/CMAKE_HIP_ARCHITECTURES /variable/CMAKE_ISPC_HEADER_DIRECTORY /variable/CMAKE_ISPC_HEADER_SUFFIX /variable/CMAKE_ISPC_INSTRUCTION_SETS 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 @@ Specify the programming language in which a source file is written. A property that can be set to indicate what programming language the source file is. If it is not set the language is determined based on the file extension. Typical values are ``CXX`` (i.e. C++), ``C``, -``CSharp``, ``CUDA``, ``Fortran``, ``ISPC``, and ``ASM``. Setting this -property for a file means this file will be compiled. Do not set this +``CSharp``, ``CUDA``, ``Fortran``, ``HIP``, ``ISPC``, and ``ASM``. Setting +this property for a file means this file will be compiled. Do not set this for headers or files that should not be compiled. .. versionchanged:: 3.20 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/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 @@ include: 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 @@ CMAKE_<LANG>_COMPILER_LAUNCHER Default value for :prop_tgt:`<LANG>_COMPILER_LAUNCHER` target property. This variable is used to initialize the property on each target as it is created. This is done only when ``<LANG>`` is ``C``, ``CXX``, ``Fortran``, -``ISPC``, ``OBJC``, ``OBJCXX``, or ``CUDA``. +``HIP``, ``ISPC``, ``OBJC``, ``OBJCXX``, or ``CUDA``. This variable is initialized to the :envvar:`CMAKE_<LANG>_COMPILER_LAUNCHER` environment variable if it is set. 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 @@ function(compiler_id_detection outvar lang) 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 @@ function(cmake_determine_compile_features lang) message(CHECK_PASS "done") + elseif(lang STREQUAL HIP AND COMMAND cmake_record_hip_compile_features) + message(CHECK_START "Detecting ${lang} compile features") + + set(CMAKE_HIP98_COMPILE_FEATURES) + set(CMAKE_HIP11_COMPILE_FEATURES) + set(CMAKE_HIP14_COMPILE_FEATURES) + set(CMAKE_HIP17_COMPILE_FEATURES) + set(CMAKE_HIP20_COMPILE_FEATURES) + set(CMAKE_HIP23_COMPILE_FEATURES) + + include("${CMAKE_ROOT}/Modules/Internal/FeatureTesting.cmake") + + cmake_record_hip_compile_features() + + if(NOT _result EQUAL 0) + message(CHECK_FAIL "failed") + return() + endif() + + if (CMAKE_HIP20_COMPILE_FEATURES AND CMAKE_HIP23_COMPILE_FEATURES) + list(REMOVE_ITEM CMAKE_HIP23_COMPILE_FEATURES ${CMAKE_HIP20_COMPILE_FEATURES}) + endif() + if (CMAKE_HIP17_COMPILE_FEATURES AND CMAKE_HIP20_COMPILE_FEATURES) + list(REMOVE_ITEM CMAKE_HIP20_COMPILE_FEATURES ${CMAKE_HIP17_COMPILE_FEATURES}) + endif() + if (CMAKE_HIP14_COMPILE_FEATURES AND CMAKE_HIP17_COMPILE_FEATURES) + list(REMOVE_ITEM CMAKE_HIP17_COMPILE_FEATURES ${CMAKE_HIP14_COMPILE_FEATURES}) + endif() + if (CMAKE_HIP11_COMPILE_FEATURES AND CMAKE_HIP14_COMPILE_FEATURES) + list(REMOVE_ITEM CMAKE_HIP14_COMPILE_FEATURES ${CMAKE_HIP11_COMPILE_FEATURES}) + endif() + if (CMAKE_HIP98_COMPILE_FEATURES AND CMAKE_HIP11_COMPILE_FEATURES) + list(REMOVE_ITEM CMAKE_HIP11_COMPILE_FEATURES ${CMAKE_HIP98_COMPILE_FEATURES}) + endif() + + if(NOT CMAKE_HIP_COMPILE_FEATURES) + set(CMAKE_HIP_COMPILE_FEATURES + ${CMAKE_HIP98_COMPILE_FEATURES} + ${CMAKE_HIP11_COMPILE_FEATURES} + ${CMAKE_HIP14_COMPILE_FEATURES} + ${CMAKE_HIP17_COMPILE_FEATURES} + ${CMAKE_HIP20_COMPILE_FEATURES} + ${CMAKE_HIP23_COMPILE_FEATURES} + ) + endif() + + set(CMAKE_HIP_COMPILE_FEATURES ${CMAKE_HIP_COMPILE_FEATURES} PARENT_SCOPE) + set(CMAKE_HIP98_COMPILE_FEATURES ${CMAKE_HIP98_COMPILE_FEATURES} PARENT_SCOPE) + set(CMAKE_HIP11_COMPILE_FEATURES ${CMAKE_HIP11_COMPILE_FEATURES} PARENT_SCOPE) + set(CMAKE_HIP14_COMPILE_FEATURES ${CMAKE_HIP14_COMPILE_FEATURES} PARENT_SCOPE) + set(CMAKE_HIP17_COMPILE_FEATURES ${CMAKE_HIP17_COMPILE_FEATURES} PARENT_SCOPE) + set(CMAKE_HIP20_COMPILE_FEATURES ${CMAKE_HIP20_COMPILE_FEATURES} PARENT_SCOPE) + set(CMAKE_HIP23_COMPILE_FEATURES ${CMAKE_HIP23_COMPILE_FEATURES} PARENT_SCOPE) + + message(CHECK_PASS "done") + endif() endfunction() 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 @@ function(CMAKE_DETERMINE_COMPILER_ID lang flagvar src) 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 @@ function(CHECK_LINKER_FLAG _lang _flag _var) 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 @@ macro(cmake_record_cuda_compile_features) unset(CMAKE_CUDA03_STANDARD__HAS_FULL_SUPPORT) endif() endmacro() + +macro(cmake_record_hip_compile_features) + set(_result 0) + if(_result EQUAL 0 AND DEFINED CMAKE_HIP23_STANDARD_COMPILE_OPTION) + _has_compiler_features_hip(23) + endif() + if(_result EQUAL 0 AND DEFINED CMAKE_HIP20_STANDARD_COMPILE_OPTION) + _has_compiler_features_hip(20) + endif() + if(_result EQUAL 0 AND DEFINED CMAKE_HIP17_STANDARD_COMPILE_OPTION) + _has_compiler_features_hip(17) + endif() + _has_compiler_features_hip(14) + _has_compiler_features_hip(11) + _has_compiler_features_hip(98) +endmacro() 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/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 @@ function(CMAKE_CHECK_COMPILER_FLAG _lang _flag _var) 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 @@ function(CMAKE_CHECK_SOURCE_COMPILES _lang _source _var) 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 @@ function(CMAKE_CHECK_SOURCE_RUNS _lang _source _var) 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 @@ macro(_record_compiler_features_cuda std) unset(lang_level_has_features) endmacro() +macro(_record_compiler_features_hip std) + list(APPEND CMAKE_HIP${std}_COMPILE_FEATURES hip_std_${std}) + + get_property(lang_level_has_features GLOBAL PROPERTY CMAKE_HIP${std}_KNOWN_FEATURES) + if(lang_level_has_features) + _record_compiler_features(HIP "${CMAKE_HIP${std}_STANDARD_COMPILE_OPTION}" CMAKE_HIP${std}_COMPILE_FEATURES) + endif() + unset(lang_level_has_features) +endmacro() + macro(_has_compiler_features lang level compile_flags feature_list) # presume all known features are supported get_property(known_features GLOBAL PROPERTY CMAKE_${lang}${level}_KNOWN_FEATURES) @@ -117,3 +127,7 @@ macro(_has_compiler_features_cuda std) list(APPEND CMAKE_CUDA${std}_COMPILE_FEATURES cuda_std_${std}) _has_compiler_features(CUDA ${std} "${CMAKE_CUDA${std}_STANDARD_COMPILE_OPTION}" CMAKE_CUDA${std}_COMPILE_FEATURES) endmacro() +macro(_has_compiler_features_hip std) + list(APPEND CMAKE_HIP${std}_COMPILE_FEATURES hip_std_${std}) + _has_compiler_features(HIP ${std} "${CMAKE_HIP${std}_STANDARD_COMPILE_OPTION}" CMAKE_HIP${std}_COMPILE_FEATURES) +endmacro() 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 @@ elseif(_WCDH_policy STREQUAL "") ) endif() +set(__skip_rocmclang TRUE) include(${CMAKE_CURRENT_LIST_DIR}/CMakeCompilerIdDetection.cmake) function(_load_compiler_variables CompilerId lang) 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 @@ void cmComputeLinkDepends::FollowLinkEntry(BFSEntry qe) // 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 @@ void cmComputeLinkDepends::AddDirectLinkEntries() 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 @@ void cmComputeLinkInformation::AddImplicitLinkInfo() this->Target->GetLinkClosure(this->Config); for (std::string const& li : lc->Languages) { - if (li == "CUDA") { + if (li == "CUDA" || li == "HIP") { // These need to go before the other implicit link information // as they could require symbols from those other library - // Currently restricted to CUDA as it is the only language + // Currently restricted as CUDA and HIP are the only languages // we have documented runtime behavior controls for this->AddRuntimeLinkLibrary(li); } 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 @@ SETUP_LANGUAGE(cuda_properties, CUDA); // NOLINTNEXTLINE(bugprone-suspicious-missing-comma) SETUP_LANGUAGE(fortran_properties, Fortran); // NOLINTNEXTLINE(bugprone-suspicious-missing-comma) +SETUP_LANGUAGE(hip_properties, HIP); +// NOLINTNEXTLINE(bugprone-suspicious-missing-comma) SETUP_LANGUAGE(objc_properties, OBJC); // NOLINTNEXTLINE(bugprone-suspicious-missing-comma) SETUP_LANGUAGE(objcxx_properties, OBJCXX); @@ -201,6 +203,8 @@ SETUP_LANGUAGE(swift_properties, Swift); std::string const kCMAKE_CUDA_ARCHITECTURES = "CMAKE_CUDA_ARCHITECTURES"; std::string const kCMAKE_CUDA_RUNTIME_LIBRARY = "CMAKE_CUDA_RUNTIME_LIBRARY"; std::string const kCMAKE_ENABLE_EXPORTS = "CMAKE_ENABLE_EXPORTS"; +std::string const kCMAKE_HIP_ARCHITECTURES = "CMAKE_HIP_ARCHITECTURES"; +std::string const kCMAKE_HIP_RUNTIME_LIBRARY = "CMAKE_HIP_RUNTIME_LIBRARY"; std::string const kCMAKE_ISPC_INSTRUCTION_SETS = "CMAKE_ISPC_INSTRUCTION_SETS"; std::string const kCMAKE_ISPC_HEADER_SUFFIX = "CMAKE_ISPC_HEADER_SUFFIX"; std::string const kCMAKE_LINK_SEARCH_END_STATIC = @@ -274,6 +278,7 @@ int cmCoreTryCompile::TryCompileCode(std::vector<std::string> const& argv, LanguageStandardState cState("C"); LanguageStandardState cudaState("CUDA"); LanguageStandardState cxxState("CXX"); + LanguageStandardState hipState("HIP"); LanguageStandardState objcState("OBJC"); LanguageStandardState objcxxState("OBJCXX"); std::vector<std::string> targets; @@ -323,6 +328,7 @@ int cmCoreTryCompile::TryCompileCode(std::vector<std::string> const& argv, } else if (cState.UpdateIfMatches(argv, i) || cxxState.UpdateIfMatches(argv, i) || cudaState.UpdateIfMatches(argv, i) || + hipState.UpdateIfMatches(argv, i) || objcState.UpdateIfMatches(argv, i) || objcxxState.UpdateIfMatches(argv, i)) { continue; @@ -428,6 +434,9 @@ int cmCoreTryCompile::TryCompileCode(std::vector<std::string> const& argv, if (!cudaState.Validate(this->Makefile)) { return -1; } + if (!hipState.Validate(this->Makefile)) { + return -1; + } if (!cxxState.Validate(this->Makefile)) { return -1; } @@ -715,6 +724,8 @@ int cmCoreTryCompile::TryCompileCode(std::vector<std::string> const& argv, vars.insert( &fortran_properties[lang_property_start], &fortran_properties[lang_property_start + lang_property_size]); + vars.insert(&hip_properties[lang_property_start], + &hip_properties[lang_property_start + lang_property_size]); vars.insert(&objc_properties[lang_property_start], &objc_properties[lang_property_start + lang_property_size]); vars.insert( @@ -727,6 +738,8 @@ int cmCoreTryCompile::TryCompileCode(std::vector<std::string> const& argv, vars.insert(kCMAKE_CUDA_ARCHITECTURES); vars.insert(kCMAKE_CUDA_RUNTIME_LIBRARY); vars.insert(kCMAKE_ENABLE_EXPORTS); + vars.insert(kCMAKE_HIP_ARCHITECTURES); + vars.insert(kCMAKE_HIP_RUNTIME_LIBRARY); vars.insert(kCMAKE_ISPC_INSTRUCTION_SETS); vars.insert(kCMAKE_ISPC_HEADER_SUFFIX); vars.insert(kCMAKE_LINK_SEARCH_END_STATIC); @@ -761,6 +774,8 @@ int cmCoreTryCompile::TryCompileCode(std::vector<std::string> const& argv, vars.insert( &fortran_properties[pie_property_start], &fortran_properties[pie_property_start + pie_property_size]); + vars.insert(&hip_properties[pie_property_start], + &hip_properties[pie_property_start + pie_property_size]); vars.insert(&objc_properties[pie_property_start], &objc_properties[pie_property_start + pie_property_size]); vars.insert( @@ -835,6 +850,7 @@ int cmCoreTryCompile::TryCompileCode(std::vector<std::string> const& argv, cState.Enabled(testLangs.find("C") != testLangs.end()); cxxState.Enabled(testLangs.find("CXX") != testLangs.end()); cudaState.Enabled(testLangs.find("CUDA") != testLangs.end()); + hipState.Enabled(testLangs.find("HIP") != testLangs.end()); objcState.Enabled(testLangs.find("OBJC") != testLangs.end()); objcxxState.Enabled(testLangs.find("OBJCXX") != testLangs.end()); @@ -842,7 +858,7 @@ int cmCoreTryCompile::TryCompileCode(std::vector<std::string> const& argv, bool honorStandard = true; if (cState.DidNone() && cxxState.DidNone() && objcState.DidNone() && - objcxxState.DidNone() && cudaState.DidNone()) { + objcxxState.DidNone() && cudaState.DidNone() && hipState.DidNone()) { switch (this->Makefile->GetPolicyStatus(cmPolicies::CMP0067)) { case cmPolicies::WARN: warnCMP0067 = this->Makefile->PolicyOptionalWarningEnabled( @@ -872,6 +888,8 @@ int cmCoreTryCompile::TryCompileCode(std::vector<std::string> const& argv, warnCMP0067, warnCMP0067Variables); cudaState.LoadUnsetPropertyValues(this->Makefile, honorStandard, warnCMP0067, warnCMP0067Variables); + hipState.LoadUnsetPropertyValues(this->Makefile, honorStandard, + warnCMP0067, warnCMP0067Variables); objcState.LoadUnsetPropertyValues(this->Makefile, honorStandard, warnCMP0067, warnCMP0067Variables); objcxxState.LoadUnsetPropertyValues(this->Makefile, honorStandard, @@ -894,6 +912,7 @@ int cmCoreTryCompile::TryCompileCode(std::vector<std::string> const& argv, cState.WriteProperties(fout, targetName); cxxState.WriteProperties(fout, targetName); cudaState.WriteProperties(fout, targetName); + hipState.WriteProperties(fout, targetName); objcState.WriteProperties(fout, targetName); objcxxState.WriteProperties(fout, targetName); 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 @@ void cmExtraCodeBlocksGenerator::CreateNewProjectFile( continue; } - // check whether it is a C/C++/CUDA implementation file + // check whether it is a C/C++/CUDA/HIP implementation file bool isCFile = false; std::string lang = s->GetOrDetermineLanguage(); - if (lang == "C" || lang == "CXX" || lang == "CUDA") { + if (lang == "C" || lang == "CXX" || lang == "CUDA" || + lang == "HIP") { std::string const& srcext = s->GetExtension(); isCFile = cm->IsACLikeSourceExtension(srcext); } 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 @@ struct CompilerIdNode : public cmGeneratorExpressionNode static const CompilerIdNode cCompilerIdNode("C"), cxxCompilerIdNode("CXX"), cudaCompilerIdNode("CUDA"), objcCompilerIdNode("OBJC"), objcxxCompilerIdNode("OBJCXX"), fortranCompilerIdNode("Fortran"), - ispcCompilerIdNode("ISPC"); + hipCompilerIdNode("HIP"), ispcCompilerIdNode("ISPC"); struct CompilerVersionNode : public cmGeneratorExpressionNode { @@ -783,7 +783,8 @@ struct CompilerVersionNode : public cmGeneratorExpressionNode static const CompilerVersionNode cCompilerVersionNode("C"), cxxCompilerVersionNode("CXX"), cudaCompilerVersionNode("CUDA"), objcCompilerVersionNode("OBJC"), objcxxCompilerVersionNode("OBJCXX"), - fortranCompilerVersionNode("Fortran"), ispcCompilerVersionNode("ISPC"); + fortranCompilerVersionNode("Fortran"), ispcCompilerVersionNode("ISPC"), + hipCompilerVersionNode("HIP"); struct PlatformIdNode : public cmGeneratorExpressionNode { @@ -2597,6 +2598,7 @@ const cmGeneratorExpressionNode* cmGeneratorExpressionNode::GetNode( { "OBJCXX_COMPILER_ID", &objcxxCompilerIdNode }, { "CUDA_COMPILER_ID", &cudaCompilerIdNode }, { "Fortran_COMPILER_ID", &fortranCompilerIdNode }, + { "HIP_COMPILER_ID", &hipCompilerIdNode }, { "VERSION_GREATER", &versionGreaterNode }, { "VERSION_GREATER_EQUAL", &versionGreaterEqNode }, { "VERSION_LESS", &versionLessNode }, @@ -2608,6 +2610,7 @@ const cmGeneratorExpressionNode* cmGeneratorExpressionNode::GetNode( { "OBJC_COMPILER_VERSION", &objcCompilerVersionNode }, { "OBJCXX_COMPILER_VERSION", &objcxxCompilerVersionNode }, { "Fortran_COMPILER_VERSION", &fortranCompilerVersionNode }, + { "HIP_COMPILER_VERSION", &hipCompilerVersionNode }, { "PLATFORM_ID", &platformIdNode }, { "COMPILE_FEATURES", &compileFeaturesNode }, { "CONFIGURATION", &configurationNode }, 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 @@ void cmGeneratorTarget::ClearSourcesCache() 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 @@ cmProp cmGeneratorTarget::GetPropertyWithPairedLanguageSupport( // Check if we should use the value set by another language. if (lang == "OBJC") { propertyValue = this->GetPropertyWithPairedLanguageSupport("C", suffix); - } else if (lang == "OBJCXX" || lang == "CUDA") { + } else if (lang == "OBJCXX" || lang == "CUDA" || lang == "HIP") { propertyValue = this->GetPropertyWithPairedLanguageSupport("CXX", suffix); } @@ -1120,7 +1121,7 @@ void cmGeneratorTarget::AppendLanguageSideEffects( std::map<std::string, std::set<cmGeneratorTarget const*>>& sideEffects) const { static const std::set<cm::string_view> LANGS_WITH_NO_SIDE_EFFECTS = { - "C"_s, "CXX"_s, "OBJC"_s, "OBJCXX"_s, "ASM"_s, "CUDA"_s, + "C"_s, "CXX"_s, "OBJC"_s, "OBJCXX"_s, "ASM"_s, "CUDA"_s, "HIP"_s }; for (auto const& lang : this->GetAllConfigCompileLanguages()) { @@ -1234,6 +1235,20 @@ bool cmGeneratorTarget::IsSystemIncludeDirectory( &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 AddLangSpecificImplicitIncludeDirectories( } } +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 @@ std::vector<BT<std::string>> cmGeneratorTarget::GetSourceFilePaths( // 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::AddISPCTargetFlags(std::string& flags) const } } +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 @@ std::vector<BT<std::string>> cmGeneratorTarget::GetIncludeDirectories( } 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 @@ std::vector<BT<std::string>> cmGeneratorTarget::GetCompileOptions( 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 @@ std::vector<BT<std::string>> cmGeneratorTarget::GetCompileFeatures( 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 @@ std::vector<BT<std::string>> cmGeneratorTarget::GetCompileDefinitions( 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 @@ std::vector<BT<std::string>> cmGeneratorTarget::GetPrecompileHeaders( 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 @@ std::vector<BT<std::string>> cmGeneratorTarget::GetLinkOptions( 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 @@ std::vector<BT<std::string>> cmGeneratorTarget::GetLinkDirectories( 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 @@ std::vector<BT<std::string>> cmGeneratorTarget::GetLinkDepends( } } 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 @@ bool cmGeneratorTarget::ComputeCompileFeatures( } // Custom updates for the CUDA standard. - if (generatorTargetLanguageStandard && language.first == "CUDA") { + if (generatorTargetLanguageStandard != nullptr && + (language.first == "CUDA")) { if (generatorTargetLanguageStandard->Value == "98") { this->LanguageStandardMap[key].Value = "03"; } @@ -6363,6 +6453,7 @@ cmLinkInterface const* cmGeneratorTarget::GetLinkInterface( iface.AllDone = true; if (iface.Exists) { this->ComputeLinkInterface(config, iface, head, secondPass); + this->ComputeLinkInterfaceRuntimeLibraries(config, iface); } } @@ -6886,6 +6977,83 @@ void cmGeneratorTarget::ComputeLinkInterfaceLibraries( } } +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 @@ const cmLinkImplementation* cmGeneratorTarget::GetLinkImplementation( 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 @@ public: void AddCUDAArchitectureFlags(std::string& flags) const; void AddCUDAToolkitFlags(std::string& flags) const; + void AddHIPArchitectureFlags(std::string& flags) const; + void AddISPCTargetFlags(std::string& flags) const; std::string GetFeatureSpecificLinkRuleVariable( @@ -1109,6 +1111,12 @@ private: 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 @@ struct cmLinkInterface : public cmLinkInterfaceLibraries { // 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 @@ struct cmLinkImplementation : public cmLinkImplementationLibraries { // 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 @@ bool cmLocalGenerator::ComputeTargetCompileFeatures() this->Makefile->GetGeneratorConfigs(cmMakefile::IncludeEmptyConfig); using LanguagePair = std::pair<std::string, std::string>; - std::vector<LanguagePair> pairedLanguages{ { "OBJC", "C" }, - { "OBJCXX", "CXX" }, - { "CUDA", "CXX" } }; + std::vector<LanguagePair> pairedLanguages{ + { "OBJC", "C" }, { "OBJCXX", "CXX" }, { "CUDA", "CXX" }, { "HIP", "CXX" } + }; std::set<LanguagePair> inferredEnabledLanguages; for (auto const& lang : pairedLanguages) { if (this->Makefile->GetState()->GetLanguageEnabled(lang.first)) { @@ -1404,8 +1404,8 @@ void cmLocalGenerator::GetDeviceLinkFlags( linkLineComputer->GetLinkerLanguage(target, config); if (pcli) { - // Compute the required cuda device link libraries when - // resolving cuda device symbols + // Compute the required device link libraries when + // resolving gpu lang device symbols this->OutputLinkLibraries(pcli, linkLineComputer, linkLibs, frameworkPath, linkPath); } @@ -1968,6 +1968,8 @@ void cmLocalGenerator::AddLanguageFlags(std::string& flags, compilerSimulateId = this->Makefile->GetSafeDefinition("CMAKE_CXX_SIMULATE_ID"); } + } else if (lang == "HIP") { + target->AddHIPArchitectureFlags(flags); } // Add VFS Overlay for Clang compilers 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 @@ void cmLocalUnixMakefileGenerator3::WriteLocalMakefile() for (LocalObjectEntry const& entry : localObjectFile.second) { if (entry.Language == "C" || entry.Language == "CXX" || entry.Language == "CUDA" || entry.Language == "Fortran" || - entry.Language == "ISPC") { - // Right now, C, C++, Fortran and CUDA have both a preprocessor and the - // ability to generate assembly code + entry.Language == "HIP" || entry.Language == "ISPC") { + // Right now, C, C++, CUDA, Fortran, HIP and ISPC have both a + // preprocessor and the ability to generate assembly code lang_has_preprocessor = true; lang_has_assembly = true; break; @@ -1550,7 +1550,7 @@ bool cmLocalUnixMakefileGenerator3::ScanDependencies( std::unique_ptr<cmDepends> scanner; if (lang == "C" || lang == "CXX" || lang == "RC" || lang == "ASM" || lang == "OBJC" || lang == "OBJCXX" || lang == "CUDA" || - lang == "ISPC") { + lang == "HIP" || lang == "ISPC") { // TODO: Handle RC (resource files) dependencies correctly. scanner = cm::make_unique<cmDependsC>(this, targetDir, lang, &validDeps); } 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 @@ void cmMakefileTargetGenerator::WriteObjectRuleFiles( std::string compilerLauncher; if (!compileCommands.empty() && (lang == "C" || lang == "CXX" || lang == "Fortran" || lang == "CUDA" || - lang == "ISPC" || lang == "OBJC" || lang == "OBJCXX")) { + lang == "HIP" || lang == "ISPC" || lang == "OBJC" || + lang == "OBJCXX")) { std::string const clauncher_prop = lang + "_COMPILER_LAUNCHER"; cmProp clauncher = this->GeneratorTarget->GetProperty(clauncher_prop); if (cmNonempty(clauncher)) { 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 @@ void cmNinjaTargetGenerator::WriteCompileRule(const std::string& lang, std::string compilerLauncher; if (!compileCmds.empty() && (lang == "C" || lang == "CXX" || lang == "Fortran" || lang == "CUDA" || - lang == "ISPC" || lang == "OBJC" || lang == "OBJCXX")) { + lang == "HIP" || lang == "ISPC" || lang == "OBJC" || + lang == "OBJCXX")) { std::string const clauncher_prop = cmStrCat(lang, "_COMPILER_LAUNCHER"); cmProp clauncher = this->GeneratorTarget->GetProperty(clauncher_prop); if (cmNonempty(clauncher)) { 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 CXX_FEATURES[] = { nullptr FOR_EACH_CXX_FEATURE( const char* const CUDA_FEATURES[] = { nullptr FOR_EACH_CUDA_FEATURE( FEATURE_STRING) }; + +const char* const HIP_FEATURES[] = { nullptr FOR_EACH_HIP_FEATURE( + FEATURE_STRING) }; #undef FEATURE_STRING struct StandardNeeded @@ -306,8 +309,7 @@ struct StanardLevelComputer }; std::unordered_map<std::string, StanardLevelComputer> StandardComputerMapping = - { - { "C", + { { "C", StanardLevelComputer{ "C", std::vector<int>{ 90, 99, 11, 17, 23 }, std::vector<std::string>{ "90", "99", "11", "17", "23" } } }, @@ -326,7 +328,10 @@ std::unordered_map<std::string, StanardLevelComputer> StandardComputerMapping = StanardLevelComputer{ "OBJCXX", std::vector<int>{ 98, 11, 14, 17, 20, 23 }, std::vector<std::string>{ "98", "11", "14", "17", "20", "23" } } }, - }; + { "HIP", + StanardLevelComputer{ + "HIP", std::vector<int>{ 98, 11, 14, 17, 20, 23 }, + std::vector<std::string>{ "98", "11", "14", "17", "20", "23" } } } }; } std::string cmStandardLevelResolver::GetCompileOptionDef( @@ -434,6 +439,13 @@ bool cmStandardLevelResolver::CompileFeatureKnown( lang = "CUDA"; return true; } + bool isHIPFeature = + std::find_if(cm::cbegin(HIP_FEATURES) + 1, cm::cend(HIP_FEATURES), + cmStrCmp(feature)) != cm::cend(HIP_FEATURES); + if (isHIPFeature) { + lang = "HIP"; + return true; + } std::ostringstream e; if (error) { e << "specified"; 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 @@ cmTarget::cmTarget(std::string const& name, cmStateEnums::TargetType type, SETUP_COMMON_LANGUAGE_PROPERTIES(CXX); SETUP_COMMON_LANGUAGE_PROPERTIES(OBJCXX); SETUP_COMMON_LANGUAGE_PROPERTIES(CUDA); + SETUP_COMMON_LANGUAGE_PROPERTIES(HIP); initProp("ANDROID_API"); initProp("ANDROID_API_MIN"); @@ -365,6 +366,8 @@ cmTarget::cmTarget(std::string const& name, cmStateEnums::TargetType type, initProp("CUDA_RESOLVE_DEVICE_SYMBOLS"); initProp("CUDA_RUNTIME_LIBRARY"); initProp("CUDA_ARCHITECTURES"); + initProp("HIP_RUNTIME_LIBRARY"); + initProp("HIP_ARCHITECTURES"); initProp("VISIBILITY_INLINES_HIDDEN"); initProp("JOB_POOL_COMPILE"); initProp("JOB_POOL_LINK"); @@ -1168,6 +1171,7 @@ void cmTarget::SetProperty(const std::string& prop, const char* value) MAKE_STATIC_PROP(C_STANDARD); MAKE_STATIC_PROP(CXX_STANDARD); MAKE_STATIC_PROP(CUDA_STANDARD); + MAKE_STATIC_PROP(HIP_STANDARD); MAKE_STATIC_PROP(OBJC_STANDARD); MAKE_STATIC_PROP(OBJCXX_STANDARD); MAKE_STATIC_PROP(COMPILE_DEFINITIONS); @@ -1354,8 +1358,8 @@ void cmTarget::SetProperty(const std::string& prop, const char* value) this->SetProperty("COMPILE_PDB_NAME", cmToCStr(tmp)); this->AddUtility(reusedFrom, false, this->impl->Makefile); } else if (prop == propC_STANDARD || prop == propCXX_STANDARD || - prop == propCUDA_STANDARD || prop == propOBJC_STANDARD || - prop == propOBJCXX_STANDARD) { + prop == propCUDA_STANDARD || prop == propHIP_STANDARD || + prop == propOBJC_STANDARD || prop == propOBJCXX_STANDARD) { if (value) { this->impl->LanguageStandardProperties[prop] = BTs<std::string>(value, this->impl->Makefile->GetBacktrace()); @@ -1461,8 +1465,8 @@ void cmTarget::AppendProperty(const std::string& prop, this->impl->Makefile->IssueMessage( MessageType::FATAL_ERROR, prop + " property may not be APPENDed."); } else if (prop == "C_STANDARD" || prop == "CXX_STANDARD" || - prop == "CUDA_STANDARD" || prop == "OBJC_STANDARD" || - prop == "OBJCXX_STANDARD") { + prop == "CUDA_STANDARD" || prop == "HIP_STANDARD" || + prop == "OBJC_STANDARD" || prop == "OBJCXX_STANDARD") { this->impl->Makefile->IssueMessage( MessageType::FATAL_ERROR, prop + " property may not be appended."); } else { diff --git a/Source/cmake.cxx b/Source/cmake.cxx index 9a866ab..eb4c0b6 100644 --- a/Source/cmake.cxx +++ b/Source/cmake.cxx @@ -218,6 +218,7 @@ cmake::cmake(Role role, cmState::Mode mode) setupExts(this->CudaFileExtensions, { "cu" }); setupExts(this->FortranFileExtensions, { "f", "F", "for", "f77", "f90", "f95", "f03" }); + setupExts(this->HipFileExtensions, { "hip" }); setupExts(this->ISPCFileExtensions, { "ispc" }); } } @@ -2462,6 +2463,8 @@ std::vector<std::string> cmake::GetAllExtensions() const // cuda extensions are also in SourceFileExtensions so we ignore it here allExt.insert(allExt.end(), this->FortranFileExtensions.ordered.begin(), this->FortranFileExtensions.ordered.end()); + allExt.insert(allExt.end(), this->HipFileExtensions.ordered.begin(), + this->HipFileExtensions.ordered.end()); allExt.insert(allExt.end(), this->ISPCFileExtensions.ordered.begin(), this->ISPCFileExtensions.ordered.end()); return allExt; 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 @@ public: return this->CLikeSourceFileExtensions.Test(ext) || this->CudaFileExtensions.Test(ext) || this->FortranFileExtensions.Test(ext) || - this->ISPCFileExtensions.Test(ext); + this->HipFileExtensions.Test(ext) || this->ISPCFileExtensions.Test(ext); } bool IsACLikeSourceExtension(cm::string_view ext) const @@ -662,6 +662,7 @@ private: FileExtensions CudaFileExtensions; FileExtensions ISPCFileExtensions; FileExtensions FortranFileExtensions; + FileExtensions HipFileExtensions; bool ClearBuildSystem = false; bool DebugTryCompile = false; bool RegenerateDuringBuild = false; @@ -838,3 +839,11 @@ private: F(cuda_std_17) \ F(cuda_std_20) \ F(cuda_std_23) + +#define FOR_EACH_HIP_FEATURE(F) \ + F(hip_std_98) \ + F(hip_std_11) \ + F(hip_std_14) \ + F(hip_std_17) \ + F(hip_std_20) \ + F(hip_std_23) 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 @@ if(BUILD_TESTING) add_subdirectory(CudaOnly) endif() + if(CMake_TEST_HIP) + add_subdirectory(HIP) + endif() + if(CMake_TEST_ISPC) add_subdirectory(ISPC) endif() diff --git a/Tests/HIP/ArchitectureOff/CMakeLists.txt b/Tests/HIP/ArchitectureOff/CMakeLists.txt new file mode 100644 index 0000000..bccb3b4 --- /dev/null +++ b/Tests/HIP/ArchitectureOff/CMakeLists.txt @@ -0,0 +1,8 @@ +cmake_minimum_required(VERSION 3.18) +project(HIPArchitecture HIP) + +# Make sure CMake doesn't pass architectures if HIP_ARCHITECTURES is OFF. +string(APPEND CMAKE_HIP_FLAGS " --offload-arch=gfx908") + +add_executable(HIPOnlyArchitectureOff main.hip) +set_property(TARGET HIPOnlyArchitectureOff PROPERTY HIP_ARCHITECTURES OFF) diff --git a/Tests/HIP/ArchitectureOff/main.hip b/Tests/HIP/ArchitectureOff/main.hip new file mode 100644 index 0000000..9256318 --- /dev/null +++ b/Tests/HIP/ArchitectureOff/main.hip @@ -0,0 +1,9 @@ +#ifdef __HIP_DEVICE_COMPILE__ +# ifndef __gfx908__ +# error "Passed architecture gfx908, but got something else." +# endif +#endif + +int main() +{ +} diff --git a/Tests/HIP/CMakeLists.txt b/Tests/HIP/CMakeLists.txt new file mode 100644 index 0000000..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..27bb929 100644 --- a/Tests/RunCMake/CMakeLists.txt +++ b/Tests/RunCMake/CMakeLists.txt @@ -452,6 +452,7 @@ function(add_RunCMake_test_try_compile) 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_include_directories) 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.CheckCompilerFlag 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("${CMAKE_GENERATOR}" MATCHES "Make|Ninja") if(DEFINED CMake_TEST_CUDA) list(APPEND CompilerLauncher_ARGS -DCMake_TEST_CUDA=${CMake_TEST_CUDA}) endif() + if(DEFINED CMake_TEST_HIP) + list(APPEND CompilerLauncher_ARGS -DCMake_TEST_HIP=${CMake_TEST_HIP}) + endif() if(DEFINED CMake_TEST_ISPC) list(APPEND CompilerLauncher_ARGS -DCMake_TEST_ISPC=${CMake_TEST_ISPC}) endif() @@ -739,7 +752,7 @@ if("${CMAKE_GENERATOR}" MATCHES "Make|Ninja") endif() add_RunCMake_test(CompilerLauncher) set_property(TEST RunCMake.CompilerLauncher APPEND - PROPERTY LABELS "CUDA;ISPC") + PROPERTY LABELS "CUDA;HIP;ISPC") add_RunCMake_test(ctest_labels_for_subprojects) add_RunCMake_test(CompilerArgs) add_RunCMake_test(LinkerLauncher) diff --git a/Tests/RunCMake/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 @@ endif() 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 @@ endif() 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 @@ endif() 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 @@ endif() 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 @@ endif() if(CMake_TEST_Fortran) list(APPEND langs Fortran) endif() +if(CMake_TEST_HIP) + list(APPEND langs HIP) +endif() if(CMake_TEST_ISPC) list(APPEND langs ISPC) endif() diff --git a/Tests/RunCMake/CompilerLauncher/main.hip b/Tests/RunCMake/CompilerLauncher/main.hip new file mode 100644 index 0000000..f8b643a --- /dev/null +++ b/Tests/RunCMake/CompilerLauncher/main.hip @@ -0,0 +1,4 @@ +int main() +{ + return 0; +} |