From b50bfc89131e55685aa41146254b37d26049b4d5 Mon Sep 17 00:00:00 2001 From: Robert Maynard Date: Fri, 28 Aug 2020 15:03:39 -0400 Subject: HIP: Add language to CMake --- Help/command/enable_language.rst | 5 +- Help/command/project.rst | 2 +- Help/envvar/HIPCXX.rst | 13 ++ Help/envvar/HIPFLAGS.rst | 15 +++ Help/manual/cmake-env-variables.7.rst | 2 + Help/manual/cmake-generator-expressions.7.rst | 22 ++++ Help/manual/cmake-properties.7.rst | 1 + Help/manual/cmake-variables.7.rst | 1 + Help/prop_sf/LANGUAGE.rst | 4 +- Help/prop_tgt/HIP_ARCHITECTURES.rst | 26 ++++ Help/prop_tgt/LANG_COMPILER_LAUNCHER.rst | 2 +- Help/release/dev/hip.rst | 5 + Help/variable/CMAKE_HIP_ARCHITECTURES.rst | 11 ++ Help/variable/CMAKE_LANG_COMPILER_LAUNCHER.rst | 2 +- Modules/CMakeDetermineCompileFeatures.cmake | 56 +++++++++ Modules/CMakeDetermineHIPCompiler.cmake | 96 +++++++++++++++ Modules/CMakeHIPCompiler.cmake.in | 59 +++++++++ Modules/CMakeHIPCompilerABI.hip | 16 +++ Modules/CMakeHIPCompilerId.hip.in | 54 +++++++++ Modules/CMakeHIPInformation.cmake | 132 +++++++++++++++++++++ Modules/CMakeTestHIPCompiler.cmake | 98 +++++++++++++++ Modules/Compiler/CMakeCommonCompilerMacros.cmake | 16 +++ Modules/Compiler/Clang-HIP.cmake | 16 +++ Modules/Compiler/ROCMClang-HIP.cmake | 45 +++++++ Modules/Internal/FeatureTesting.cmake | 14 +++ Source/cmComputeLinkInformation.cxx | 4 +- Source/cmCoreTryCompile.cxx | 21 +++- Source/cmExtraCodeBlocksGenerator.cxx | 5 +- Source/cmGeneratorExpressionNode.cxx | 7 +- Source/cmGeneratorTarget.cxx | 24 +++- Source/cmGeneratorTarget.h | 2 + Source/cmLocalGenerator.cxx | 12 +- Source/cmLocalUnixMakefileGenerator3.cxx | 8 +- Source/cmMakefileTargetGenerator.cxx | 3 +- Source/cmNinjaTargetGenerator.cxx | 3 +- Source/cmStandardLevelResolver.cxx | 18 ++- Source/cmTarget.cxx | 12 +- Source/cmake.cxx | 3 + Source/cmake.h | 11 +- Tests/CMakeLists.txt | 4 + Tests/HIP/ArchitectureOff/CMakeLists.txt | 8 ++ Tests/HIP/ArchitectureOff/main.hip | 9 ++ Tests/HIP/CMakeLists.txt | 12 ++ Tests/HIP/CompileFlags/CMakeLists.txt | 8 ++ Tests/HIP/CompileFlags/main.hip | 13 ++ Tests/HIP/EnableStandard/CMakeLists.txt | 20 ++++ Tests/HIP/EnableStandard/main.hip | 23 ++++ Tests/HIP/EnableStandard/shared.hip | 15 +++ Tests/HIP/EnableStandard/static.cxx | 9 ++ Tests/HIP/MathFunctions/CMakeLists.txt | 18 +++ Tests/HIP/MathFunctions/main.hip | 40 +++++++ Tests/HIP/TryCompile/CMakeLists.txt | 14 +++ Tests/HIP/TryCompile/device_function.hip | 17 +++ Tests/HIP/TryCompile/main.hip | 8 ++ Tests/HIP/WithDefs/CMakeLists.txt | 36 ++++++ Tests/HIP/WithDefs/inc_hip/inc_hip.h | 1 + Tests/HIP/WithDefs/main.hip.cpp | 89 ++++++++++++++ Tests/RunCMake/CMakeLists.txt | 5 +- Tests/RunCMake/CompilerLauncher/HIP-common.cmake | 5 + .../CompilerLauncher/HIP-env-Build-stdout.txt | 1 + .../HIP-env-launch-Build-stdout.txt | 1 + Tests/RunCMake/CompilerLauncher/HIP-env.cmake | 1 + .../CompilerLauncher/HIP-launch-Build-stdout.txt | 1 + .../RunCMake/CompilerLauncher/HIP-launch-env.cmake | 3 + Tests/RunCMake/CompilerLauncher/HIP-launch.cmake | 3 + Tests/RunCMake/CompilerLauncher/HIP.cmake | 2 + Tests/RunCMake/CompilerLauncher/RunCMakeTest.cmake | 3 + Tests/RunCMake/CompilerLauncher/main.hip | 4 + 68 files changed, 1183 insertions(+), 36 deletions(-) create mode 100644 Help/envvar/HIPCXX.rst create mode 100644 Help/envvar/HIPFLAGS.rst create mode 100644 Help/prop_tgt/HIP_ARCHITECTURES.rst create mode 100644 Help/release/dev/hip.rst create mode 100644 Help/variable/CMAKE_HIP_ARCHITECTURES.rst create mode 100644 Modules/CMakeDetermineHIPCompiler.cmake create mode 100644 Modules/CMakeHIPCompiler.cmake.in create mode 100644 Modules/CMakeHIPCompilerABI.hip create mode 100644 Modules/CMakeHIPCompilerId.hip.in create mode 100644 Modules/CMakeHIPInformation.cmake create mode 100644 Modules/CMakeTestHIPCompiler.cmake create mode 100644 Modules/Compiler/Clang-HIP.cmake create mode 100644 Modules/Compiler/ROCMClang-HIP.cmake create mode 100644 Tests/HIP/ArchitectureOff/CMakeLists.txt create mode 100644 Tests/HIP/ArchitectureOff/main.hip create mode 100644 Tests/HIP/CMakeLists.txt create mode 100644 Tests/HIP/CompileFlags/CMakeLists.txt create mode 100644 Tests/HIP/CompileFlags/main.hip create mode 100644 Tests/HIP/EnableStandard/CMakeLists.txt create mode 100644 Tests/HIP/EnableStandard/main.hip create mode 100644 Tests/HIP/EnableStandard/shared.hip create mode 100644 Tests/HIP/EnableStandard/static.cxx create mode 100644 Tests/HIP/MathFunctions/CMakeLists.txt create mode 100644 Tests/HIP/MathFunctions/main.hip create mode 100644 Tests/HIP/TryCompile/CMakeLists.txt create mode 100644 Tests/HIP/TryCompile/device_function.hip create mode 100644 Tests/HIP/TryCompile/main.hip create mode 100644 Tests/HIP/WithDefs/CMakeLists.txt create mode 100644 Tests/HIP/WithDefs/inc_hip/inc_hip.h create mode 100644 Tests/HIP/WithDefs/main.hip.cpp create mode 100644 Tests/RunCMake/CompilerLauncher/HIP-common.cmake create mode 100644 Tests/RunCMake/CompilerLauncher/HIP-env-Build-stdout.txt create mode 100644 Tests/RunCMake/CompilerLauncher/HIP-env-launch-Build-stdout.txt create mode 100644 Tests/RunCMake/CompilerLauncher/HIP-env.cmake create mode 100644 Tests/RunCMake/CompilerLauncher/HIP-launch-Build-stdout.txt create mode 100644 Tests/RunCMake/CompilerLauncher/HIP-launch-env.cmake create mode 100644 Tests/RunCMake/CompilerLauncher/HIP-launch.cmake create mode 100644 Tests/RunCMake/CompilerLauncher/HIP.cmake create mode 100644 Tests/RunCMake/CompilerLauncher/main.hip 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 _COMPILER>`. For any configuration +run (including the first), the environment variable will be ignored if the +:variable:`CMAKE_HIP_COMPILER _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 _FLAGS>`. For any configuration +run (including the first), the environment variable will be ignored if +the :variable:`CMAKE_HIP_FLAGS _FLAGS>` variable is defined. + +See also :variable:`CMAKE_HIP_FLAGS_INIT _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 775067a..ee73f18 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__COMPILER_ID` variable. +.. genex:: $ + + 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__COMPILER_ID` variable. + .. genex:: $ 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__COMPILER_VERSION` variable. +.. genex:: $ + + ``1`` if the version of the HIP compiler matches ``version``, otherwise ``0``. + See also the :variable:`CMAKE__COMPILER_VERSION` variable. + .. genex:: $ ``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__COMPILER_ID` variable. +.. genex:: $ + + The CMake's compiler id of the HIP compiler used. + See also the :variable:`CMAKE__COMPILER_ID` variable. + .. genex:: $ 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__COMPILER_VERSION` variable. +.. genex:: $ + + The version of the HIP compiler used. + See also the :variable:`CMAKE__COMPILER_VERSION` variable. + .. genex:: $ 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..052f472 --- /dev/null +++ b/Help/prop_tgt/HIP_ARCHITECTURES.rst @@ -0,0 +1,26 @@ +HIP_ARCHITECTURES +----------------- + +.. versionadded:: 3.21 + +List of AMD GPU architectures to generate device code for. + +An empty or false value (e.g. ``OFF``) defers architecture generation to compiler +defaults. + +This property is initialized by the value of the :variable:`CMAKE_HIP_ARCHITECTURES` +variable if it is set when a target is created. + +The HIP compilation model has two modes: whole and separable. Whole compilation +generates device code at compile time. Separable compilation generates device +code at link time. Therefore the ``HIP_ARCHITECTURES`` target property should +be set on targets that compile or link with any HIP sources. + +Examples +^^^^^^^^ + +.. code-block:: cmake + + set_property(TARGET tgt PROPERTY HIP_ARCHITECTURES gfx801 gfx900) + +Generates code for both ``gfx801`` and ``gfx900``. 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 ```` is ``C``, ``CXX``, -``Fortran``, ``ISPC``, ``OBJC``, ``OBJCXX``, or ``CUDA``. +``Fortran``, ``HIP``, ``ISPC``, ``OBJC``, ``OBJCXX``, or ``CUDA``. Specify a :ref:`semicolon-separated list ` 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..333f021 --- /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 ``OFF``. + +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_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__COMPILER_LAUNCHER Default value for :prop_tgt:`_COMPILER_LAUNCHER` target property. This variable is used to initialize the property on each target as it is created. This is done only when ```` is ``C``, ``CXX``, ``Fortran``, -``ISPC``, ``OBJC``, ``OBJCXX``, or ``CUDA``. +``HIP``, ``ISPC``, ``OBJC``, ``OBJCXX``, or ``CUDA``. This variable is initialized to the :envvar:`CMAKE__COMPILER_LAUNCHER` environment variable if it is set. 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/CMakeDetermineHIPCompiler.cmake b/Modules/CMakeDetermineHIPCompiler.cmake new file mode 100644 index 0000000..87885c0 --- /dev/null +++ b/Modules/CMakeDetermineHIPCompiler.cmake @@ -0,0 +1,96 @@ +# Distributed under the OSI-approved BSD 3-Clause License. See accompanying +# file Copyright.txt or https://cmake.org/licensing for details. + +include(${CMAKE_ROOT}/Modules/CMakeDetermineCompiler.cmake) +include(${CMAKE_ROOT}/Modules/CMakeParseImplicitLinkInfo.cmake) + +if( NOT ( ("${CMAKE_GENERATOR}" MATCHES "Make") OR + ("${CMAKE_GENERATOR}" MATCHES "Ninja") ) ) + message(FATAL_ERROR "HIP language not currently supported by \"${CMAKE_GENERATOR}\" generator") +endif() + + +if(NOT CMAKE_HIP_COMPILER) + set(CMAKE_HIP_COMPILER_INIT NOTFOUND) + + # prefer the environment variable HIPCXX + if(NOT $ENV{HIPCXX} STREQUAL "") + get_filename_component(CMAKE_HIP_COMPILER_INIT $ENV{HIPCXX} PROGRAM PROGRAM_ARGS CMAKE_HIP_FLAGS_ENV_INIT) + if(CMAKE_HIP_FLAGS_ENV_INIT) + set(CMAKE_HIP_COMPILER_ARG1 "${CMAKE_HIP_FLAGS_ENV_INIT}" CACHE STRING "Arguments to CXX compiler") + endif() + if(NOT EXISTS ${CMAKE_HIP_COMPILER_INIT}) + message(FATAL_ERROR "Could not find compiler set in environment variable HIPCXX:\n$ENV{HIPCXX}.\n${CMAKE_HIP_COMPILER_INIT}") + endif() + endif() + + # finally list compilers to try + if(NOT CMAKE_HIP_COMPILER_INIT) + set(CMAKE_HIP_COMPILER_LIST hipcc clang++) + endif() + + _cmake_find_compiler(HIP) +else() + _cmake_find_compiler_path(HIP) +endif() + +mark_as_advanced(CMAKE_HIP_COMPILER) + +# Build a small source file to identify the compiler. +if(NOT CMAKE_HIP_COMPILER_ID_RUN) + set(CMAKE_HIP_COMPILER_ID_RUN 1) + + # Try to identify the compiler. + set(CMAKE_HIP_COMPILER_ID) + set(CMAKE_HIP_PLATFORM_ID) + file(READ ${CMAKE_ROOT}/Modules/CMakePlatformId.h.in + CMAKE_HIP_COMPILER_ID_PLATFORM_CONTENT) + + list(APPEND CMAKE_HIP_COMPILER_ID_TEST_FLAGS_FIRST "-v") + + include(${CMAKE_ROOT}/Modules/CMakeDetermineCompilerId.cmake) + CMAKE_DETERMINE_COMPILER_ID(HIP HIPFLAGS CMakeHIPCompilerId.hip) + + _cmake_find_compiler_sysroot(HIP) + +endif() + +if (NOT _CMAKE_TOOLCHAIN_LOCATION) + get_filename_component(_CMAKE_TOOLCHAIN_LOCATION "${CMAKE_HIP_COMPILER}" PATH) +endif () + +set(_CMAKE_PROCESSING_LANGUAGE "HIP") +include(CMakeFindBinUtils) +include(Compiler/${CMAKE_HIP_COMPILER_ID}-FindBinUtils OPTIONAL) +unset(_CMAKE_PROCESSING_LANGUAGE) + +if(CMAKE_HIP_COMPILER_SYSROOT) + string(CONCAT _SET_CMAKE_HIP_COMPILER_SYSROOT + "set(CMAKE_HIP_COMPILER_SYSROOT \"${CMAKE_HIP_COMPILER_SYSROOT}\")\n" + "set(CMAKE_COMPILER_SYSROOT \"${CMAKE_HIP_COMPILER_SYSROOT}\")") +else() + set(_SET_CMAKE_HIP_COMPILER_SYSROOT "") +endif() + +if(CMAKE_HIP_COMPILER_ARCHITECTURE_ID) + set(_SET_CMAKE_HIP_COMPILER_ARCHITECTURE_ID + "set(CMAKE_HIP_COMPILER_ARCHITECTURE_ID ${CMAKE_HIP_COMPILER_ARCHITECTURE_ID})") +else() + set(_SET_CMAKE_HIP_COMPILER_ARCHITECTURE_ID "") +endif() + +if(MSVC_HIP_ARCHITECTURE_ID) + set(SET_MSVC_HIP_ARCHITECTURE_ID + "set(MSVC_HIP_ARCHITECTURE_ID ${MSVC_HIP_ARCHITECTURE_ID})") +endif() + +if(NOT DEFINED CMAKE_HIP_ARCHITECTURES) + set(CMAKE_HIP_ARCHITECTURES "OFF" CACHE STRING "HIP architectures") +endif() + +# configure variables set in this file for fast reload later on +configure_file(${CMAKE_ROOT}/Modules/CMakeHIPCompiler.cmake.in + ${CMAKE_PLATFORM_INFO_DIR}/CMakeHIPCompiler.cmake + @ONLY + ) +set(CMAKE_HIP_COMPILER_ENV_VAR "HIPCXX") diff --git a/Modules/CMakeHIPCompiler.cmake.in b/Modules/CMakeHIPCompiler.cmake.in new file mode 100644 index 0000000..eeb51b2 --- /dev/null +++ b/Modules/CMakeHIPCompiler.cmake.in @@ -0,0 +1,59 @@ +set(CMAKE_HIP_COMPILER "@CMAKE_HIP_COMPILER@") +set(CMAKE_HIP_COMPILER_ID "@CMAKE_HIP_COMPILER_ID@") +set(CMAKE_HIP_COMPILER_VERSION "@CMAKE_HIP_COMPILER_VERSION@") +set(CMAKE_HIP_STANDARD_COMPUTED_DEFAULT "@CMAKE_HIP_STANDARD_COMPUTED_DEFAULT@") +set(CMAKE_HIP_COMPILE_FEATURES "@CMAKE_HIP_COMPILE_FEATURES@") +set(CMAKE_HIP98_COMPILE_FEATURES "@CMAKE_HIP03_COMPILE_FEATURES@") +set(CMAKE_HIP11_COMPILE_FEATURES "@CMAKE_HIP11_COMPILE_FEATURES@") +set(CMAKE_HIP14_COMPILE_FEATURES "@CMAKE_HIP14_COMPILE_FEATURES@") +set(CMAKE_HIP17_COMPILE_FEATURES "@CMAKE_HIP17_COMPILE_FEATURES@") +set(CMAKE_HIP20_COMPILE_FEATURES "@CMAKE_HIP20_COMPILE_FEATURES@") +set(CMAKE_HIP23_COMPILE_FEATURES "@CMAKE_HIP23_COMPILE_FEATURES@") + +set(CMAKE_HIP_PLATFORM_ID "@CMAKE_HIP_PLATFORM_ID@") +set(CMAKE_HIP_SIMULATE_ID "@CMAKE_HIP_SIMULATE_ID@") +set(CMAKE_HIP_COMPILER_FRONTEND_VARIANT "@CMAKE_HIP_COMPILER_FRONTEND_VARIANT@") +set(CMAKE_HIP_SIMULATE_VERSION "@CMAKE_HIP_SIMULATE_VERSION@") +@SET_MSVC_HIP_ARCHITECTURE_ID@ +@_SET_CMAKE_HIP_COMPILER_SYSROOT@ + +set(CMAKE_HIP_COMPILER_ENV_VAR "HIPCXX") + +set(CMAKE_HIP_COMPILER_LOADED 1) +set(CMAKE_HIP_COMPILER_ID_RUN 1) +set(CMAKE_HIP_SOURCE_FILE_EXTENSIONS hip) +set(CMAKE_HIP_LINKER_PREFERENCE 15) +set(CMAKE_HIP_LINKER_PREFERENCE_PROPAGATES 1) + +set(CMAKE_HIP_SIZEOF_DATA_PTR "@CMAKE_HIP_SIZEOF_DATA_PTR@") +set(CMAKE_HIP_COMPILER_ABI "@CMAKE_HIP_COMPILER_ABI@") +set(CMAKE_HIP_LIBRARY_ARCHITECTURE "@CMAKE_HIP_LIBRARY_ARCHITECTURE@") + +if(CMAKE_HIP_SIZEOF_DATA_PTR) + set(CMAKE_SIZEOF_VOID_P "${CMAKE_HIP_SIZEOF_DATA_PTR}") +endif() + +if(CMAKE_HIP_COMPILER_ABI) + set(CMAKE_INTERNAL_PLATFORM_ABI "${CMAKE_HIP_COMPILER_ABI}") +endif() + +if(CMAKE_HIP_LIBRARY_ARCHITECTURE) + set(CMAKE_LIBRARY_ARCHITECTURE "@CMAKE_HIP_LIBRARY_ARCHITECTURE@") +endif() + +set(CMAKE_HIP_TOOLKIT_INCLUDE_DIRECTORIES "@CMAKE_HIP_TOOLKIT_INCLUDE_DIRECTORIES@") + +set(CMAKE_HIP_IMPLICIT_INCLUDE_DIRECTORIES "@CMAKE_HIP_IMPLICIT_INCLUDE_DIRECTORIES@") +set(CMAKE_HIP_IMPLICIT_LINK_LIBRARIES "@CMAKE_HIP_IMPLICIT_LINK_LIBRARIES@") +set(CMAKE_HIP_IMPLICIT_LINK_DIRECTORIES "@CMAKE_HIP_IMPLICIT_LINK_DIRECTORIES@") +set(CMAKE_HIP_IMPLICIT_LINK_FRAMEWORK_DIRECTORIES "@CMAKE_HIP_IMPLICIT_LINK_FRAMEWORK_DIRECTORIES@") + +set(CMAKE_HIP_COMPILER_DEVICE_LIBRARY_ROOT_DIR "@CMAKE_HIP_COMPILER_DEVICE_LIBRARY_ROOT_DIR@") +@_SET_CMAKE_HIP_DEVICE_LIBRARY_INCLUSION@ + +set(CMAKE_AR "@CMAKE_AR@") +set(CMAKE_HIP_COMPILER_AR "@CMAKE_HIP_COMPILER_AR@") +set(CMAKE_RANLIB "@CMAKE_RANLIB@") +set(CMAKE_HIP_COMPILER_RANLIB "@CMAKE_HIP_COMPILER_RANLIB@") +set(CMAKE_LINKER "@CMAKE_LINKER@") +set(CMAKE_MT "@CMAKE_MT@") 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..df1d98a --- /dev/null +++ b/Modules/CMakeHIPInformation.cmake @@ -0,0 +1,132 @@ +# Distributed under the OSI-approved BSD 3-Clause License. See accompanying +# file Copyright.txt or https://cmake.org/licensing for details. + +if(UNIX) + set(CMAKE_HIP_OUTPUT_EXTENSION .o) +else() + set(CMAKE_HIP_OUTPUT_EXTENSION .obj) +endif() +set(CMAKE_INCLUDE_FLAG_HIP "-I") + +# Load compiler-specific information. +if(CMAKE_HIP_COMPILER_ID) + include(Compiler/${CMAKE_HIP_COMPILER_ID}-HIP OPTIONAL) +endif() + +# load the system- and compiler specific files +if(CMAKE_HIP_COMPILER_ID) + # load a hardware specific file, mostly useful for embedded compilers + if(CMAKE_SYSTEM_PROCESSOR) + include(Platform/${CMAKE_EFFECTIVE_SYSTEM_NAME}-${CMAKE_HIP_COMPILER_ID}-HIP-${CMAKE_SYSTEM_PROCESSOR} OPTIONAL) + endif() + include(Platform/${CMAKE_EFFECTIVE_SYSTEM_NAME}-${CMAKE_HIP_COMPILER_ID}-HIP OPTIONAL) +endif() + + +if(NOT CMAKE_SHARED_LIBRARY_RUNTIME_HIP_FLAG) + set(CMAKE_SHARED_LIBRARY_RUNTIME_HIP_FLAG ${CMAKE_SHARED_LIBRARY_RUNTIME_C_FLAG}) +endif() + +if(NOT CMAKE_SHARED_LIBRARY_RUNTIME_HIP_FLAG_SEP) + set(CMAKE_SHARED_LIBRARY_RUNTIME_HIP_FLAG_SEP ${CMAKE_SHARED_LIBRARY_RUNTIME_C_FLAG_SEP}) +endif() + +if(NOT CMAKE_SHARED_LIBRARY_RPATH_LINK_HIP_FLAG) + set(CMAKE_SHARED_LIBRARY_RPATH_LINK_HIP_FLAG ${CMAKE_SHARED_LIBRARY_RPATH_LINK_C_FLAG}) +endif() + +if(NOT DEFINED CMAKE_EXE_EXPORTS_HIP_FLAG) + set(CMAKE_EXE_EXPORTS_HIP_FLAG ${CMAKE_EXE_EXPORTS_C_FLAG}) +endif() + +if(NOT DEFINED CMAKE_SHARED_LIBRARY_SONAME_HIP_FLAG) + set(CMAKE_SHARED_LIBRARY_SONAME_HIP_FLAG ${CMAKE_SHARED_LIBRARY_SONAME_C_FLAG}) +endif() + +if(NOT CMAKE_EXECUTABLE_RUNTIME_HIP_FLAG) + set(CMAKE_EXECUTABLE_RUNTIME_HIP_FLAG ${CMAKE_SHARED_LIBRARY_RUNTIME_HIP_FLAG}) +endif() + +if(NOT CMAKE_EXECUTABLE_RUNTIME_HIP_FLAG_SEP) + set(CMAKE_EXECUTABLE_RUNTIME_HIP_FLAG_SEP ${CMAKE_SHARED_LIBRARY_RUNTIME_HIP_FLAG_SEP}) +endif() + +if(NOT CMAKE_EXECUTABLE_RPATH_LINK_HIP_FLAG) + set(CMAKE_EXECUTABLE_RPATH_LINK_HIP_FLAG ${CMAKE_SHARED_LIBRARY_RPATH_LINK_HIP_FLAG}) +endif() + +if(NOT DEFINED CMAKE_SHARED_LIBRARY_LINK_HIP_WITH_RUNTIME_PATH) + set(CMAKE_SHARED_LIBRARY_LINK_HIP_WITH_RUNTIME_PATH ${CMAKE_SHARED_LIBRARY_LINK_C_WITH_RUNTIME_PATH}) +endif() + + +# for most systems a module is the same as a shared library +# so unless the variable CMAKE_MODULE_EXISTS is set just +# copy the values from the LIBRARY variables +if(NOT CMAKE_MODULE_EXISTS) + set(CMAKE_SHARED_MODULE_HIP_FLAGS ${CMAKE_SHARED_LIBRARY_HIP_FLAGS}) + set(CMAKE_SHARED_MODULE_CREATE_HIP_FLAGS ${CMAKE_SHARED_LIBRARY_CREATE_HIP_FLAGS}) +endif() + +# add the flags to the cache based +# on the initial values computed in the platform/*.cmake files +# use _INIT variables so that this only happens the first time +# and you can set these flags in the cmake cache +set(CMAKE_HIP_FLAGS_INIT "$ENV{HIPFLAGS} ${CMAKE_HIP_FLAGS_INIT}") + +cmake_initialize_per_config_variable(CMAKE_HIP_FLAGS "Flags used by the HIP compiler") + +if(CMAKE_HIP_STANDARD_LIBRARIES_INIT) + set(CMAKE_HIP_STANDARD_LIBRARIES "${CMAKE_HIP_STANDARD_LIBRARIES_INIT}" + CACHE STRING "Libraries linked by default with all HIP applications.") + mark_as_advanced(CMAKE_HIP_STANDARD_LIBRARIES) +endif() + +if(NOT CMAKE_HIP_COMPILER_LAUNCHER AND DEFINED ENV{CMAKE_HIP_COMPILER_LAUNCHER}) + set(CMAKE_HIP_COMPILER_LAUNCHER "$ENV{CMAKE_HIP_COMPILER_LAUNCHER}" + CACHE STRING "Compiler launcher for HIP.") +endif() + +include(CMakeCommonLanguageInclude) + +# now define the following rules: +# CMAKE_HIP_CREATE_SHARED_LIBRARY +# CMAKE_HIP_CREATE_SHARED_MODULE +# CMAKE_HIP_COMPILE_OBJECT +# CMAKE_HIP_LINK_EXECUTABLE + +# create a shared library +if(NOT CMAKE_HIP_CREATE_SHARED_LIBRARY) + set(CMAKE_HIP_CREATE_SHARED_LIBRARY + " -o ") +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 " qc ") +endif() +if(NOT DEFINED CMAKE_HIP_ARCHIVE_APPEND) + set(CMAKE_HIP_ARCHIVE_APPEND " q ") +endif() +if(NOT DEFINED CMAKE_HIP_ARCHIVE_FINISH) + set(CMAKE_HIP_ARCHIVE_FINISH " ") +endif() + +# compile a HIP file into an object file +if(NOT CMAKE_HIP_COMPILE_OBJECT) + set(CMAKE_HIP_COMPILE_OBJECT + " -o ${_CMAKE_COMPILE_AS_HIP_FLAG} -c ") +endif() + +# compile a cu file into an executable +if(NOT CMAKE_HIP_LINK_EXECUTABLE) + set(CMAKE_HIP_LINK_EXECUTABLE + " -o ") +endif() + +set(CMAKE_HIP_INFORMATION_LOADED 1) diff --git a/Modules/CMakeTestHIPCompiler.cmake b/Modules/CMakeTestHIPCompiler.cmake new file mode 100644 index 0000000..da70f8d --- /dev/null +++ b/Modules/CMakeTestHIPCompiler.cmake @@ -0,0 +1,98 @@ +# Distributed under the OSI-approved BSD 3-Clause License. See accompanying +# file Copyright.txt or https://cmake.org/licensing for details. + + +if(CMAKE_HIP_COMPILER_FORCED) + # The compiler configuration was forced by the user. + # Assume the user has configured all compiler information. + set(CMAKE_HIP_COMPILER_WORKS TRUE) + return() +endif() + +set(__CMAKE_HIP_FLAGS "${CMAKE_HIP_FLAGS}") +string(APPEND CMAKE_HIP_FLAGS "--cuda-host-only") + +include(CMakeTestCompilerCommon) + +# work around enforced code signing and / or missing executable target type +set(__CMAKE_SAVED_TRY_COMPILE_TARGET_TYPE ${CMAKE_TRY_COMPILE_TARGET_TYPE}) +if(_CMAKE_FEATURE_DETECTION_TARGET_TYPE) + set(CMAKE_TRY_COMPILE_TARGET_TYPE ${_CMAKE_FEATURE_DETECTION_TARGET_TYPE}) +endif() + +# Remove any cached result from an older CMake version. +# We now store this in CMakeHIPCompiler.cmake. +unset(CMAKE_HIP_COMPILER_WORKS CACHE) + +# Try to identify the ABI and configure it into CMakeHIPCompiler.cmake +include(${CMAKE_ROOT}/Modules/CMakeDetermineCompilerABI.cmake) +CMAKE_DETERMINE_COMPILER_ABI(HIP ${CMAKE_ROOT}/Modules/CMakeHIPCompilerABI.hip) +if(CMAKE_HIP_ABI_COMPILED) + # The compiler worked so skip dedicated test below. + set(CMAKE_HIP_COMPILER_WORKS TRUE) + message(STATUS "Check for working HIP compiler: ${CMAKE_HIP_COMPILER} - skipped") +endif() + +# This file is used by EnableLanguage in cmGlobalGenerator to +# determine that the selected C++ compiler can actually compile +# and link the most basic of programs. If not, a fatal error +# is set and cmake stops processing commands and will not generate +# any makefiles or projects. +if(NOT CMAKE_HIP_COMPILER_WORKS) + PrintTestCompilerStatus("HIP") + __TestCompiler_setTryCompileTargetType() + file(WRITE ${CMAKE_BINARY_DIR}${CMAKE_FILES_DIRECTORY}/CMakeTmp/testHIPCompiler.hip + "#ifndef __HIP__\n" + "# error \"The CMAKE_HIP_COMPILER is set to a C/CXX compiler\"\n" + "#endif\n" + "int main(){return 0;}\n") + try_compile(CMAKE_HIP_COMPILER_WORKS ${CMAKE_BINARY_DIR} + ${CMAKE_BINARY_DIR}${CMAKE_FILES_DIRECTORY}/CMakeTmp/testHIPCompiler.hip + OUTPUT_VARIABLE __CMAKE_HIP_COMPILER_OUTPUT) + # Move result from cache to normal variable. + set(CMAKE_HIP_COMPILER_WORKS ${CMAKE_HIP_COMPILER_WORKS}) + unset(CMAKE_HIP_COMPILER_WORKS CACHE) + __TestCompiler_restoreTryCompileTargetType() + if(NOT CMAKE_HIP_COMPILER_WORKS) + PrintTestCompilerResult(CHECK_FAIL "broken") + file(APPEND ${CMAKE_BINARY_DIR}${CMAKE_FILES_DIRECTORY}/CMakeError.log + "Determining if the HIP compiler works failed with " + "the following output:\n${__CMAKE_HIP_COMPILER_OUTPUT}\n\n") + string(REPLACE "\n" "\n " _output "${__CMAKE_HIP_COMPILER_OUTPUT}") + message(FATAL_ERROR "The HIP compiler\n \"${CMAKE_HIP_COMPILER}\"\n" + "is not able to compile a simple test program.\nIt fails " + "with the following output:\n ${_output}\n\n" + "CMake will not be able to correctly generate this project.") + endif() + PrintTestCompilerResult(CHECK_PASS "works") + file(APPEND ${CMAKE_BINARY_DIR}${CMAKE_FILES_DIRECTORY}/CMakeOutput.log + "Determining if the HIP compiler works passed with " + "the following output:\n${__CMAKE_HIP_COMPILER_OUTPUT}\n\n") +endif() + +set(CMAKE_HIP_FLAGS "${__CMAKE_HIP_FLAGS}") +unset(__CMAKE_HIP_FLAGS) + + +# Try to identify the compiler features +include(${CMAKE_ROOT}/Modules/CMakeDetermineCompileFeatures.cmake) +CMAKE_DETERMINE_COMPILE_FEATURES(HIP) + +# Re-configure to save learned information. +configure_file( + ${CMAKE_ROOT}/Modules/CMakeHIPCompiler.cmake.in + ${CMAKE_PLATFORM_INFO_DIR}/CMakeHIPCompiler.cmake + @ONLY + ) +include(${CMAKE_PLATFORM_INFO_DIR}/CMakeHIPCompiler.cmake) + +if(CMAKE_HIP_SIZEOF_DATA_PTR) + foreach(f ${CMAKE_HIP_ABI_FILES}) + include(${f}) + endforeach() + unset(CMAKE_HIP_ABI_FILES) +endif() + +set(CMAKE_TRY_COMPILE_TARGET_TYPE ${__CMAKE_SAVED_TRY_COMPILE_TARGET_TYPE}) +unset(__CMAKE_SAVED_TRY_COMPILE_TARGET_TYPE) +unset(__CMAKE_HIP_COMPILER_OUTPUT) 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..6ac39de --- /dev/null +++ b/Modules/Compiler/Clang-HIP.cmake @@ -0,0 +1,16 @@ +include(Compiler/Clang) +__compiler_clang(HIP) +__compiler_clang_cxx_standards(HIP) + +set(_CMAKE_COMPILE_AS_HIP_FLAG "-x hip") +set(_CMAKE_HIP_RDC_FLAG "-fgpu-rdc") + +if(NOT "x${CMAKE_HIP_SIMULATE_ID}" STREQUAL "xMSVC") + set(CMAKE_HIP_COMPILE_OPTIONS_VISIBILITY_INLINES_HIDDEN "-fvisibility-inlines-hidden") + + string(APPEND CMAKE_HIP_FLAGS_DEBUG_INIT " -O") +endif() + +set(CMAKE_HIP_RUNTIME_LIBRARY_DEFAULT "SHARED") +set(CMAKE_HIP_RUNTIME_LIBRARY_LINK_OPTIONS_STATIC "") +set(CMAKE_HIP_RUNTIME_LIBRARY_LINK_OPTIONS_SHARED "") diff --git a/Modules/Compiler/ROCMClang-HIP.cmake b/Modules/Compiler/ROCMClang-HIP.cmake new file mode 100644 index 0000000..9e792b5 --- /dev/null +++ b/Modules/Compiler/ROCMClang-HIP.cmake @@ -0,0 +1,45 @@ +include(Compiler/ROCMClang) +__compiler_rocmclang(HIP) + +set(_CMAKE_COMPILE_AS_HIP_FLAG "-x hip") +set(_CMAKE_HIP_RDC_FLAG "-fgpu-rdc") + +if(NOT "x${CMAKE_${lang}_SIMULATE_ID}" STREQUAL "xMSVC") + set(CMAKE_HIP_COMPILE_OPTIONS_VISIBILITY_INLINES_HIDDEN "-fvisibility-inlines-hidden") + string(APPEND CMAKE_HIP_FLAGS_DEBUG_INIT " -O") +endif() + +if(CMAKE_HIP_SIMULATE_ID STREQUAL "GNU") + set(CMAKE_HIP_LINKER_WRAPPER_FLAG "-Wl,") + set(CMAKE_HIP_LINKER_WRAPPER_FLAG_SEP ",") +elseif(CMAKE_HIP_SIMULATE_ID STREQUAL "Clang") + set(CMAKE_HIP_LINKER_WRAPPER_FLAG "-Xlinker" " ") + set(CMAKE_HIP_LINKER_WRAPPER_FLAG_SEP) +endif() + +if(NOT CMAKE_HIP_COMPILER_VERSION VERSION_LESS 1.0) + set(CMAKE_HIP98_STANDARD_COMPILE_OPTION "-std=c++98") + set(CMAKE_HIP98_EXTENSION_COMPILE_OPTION "-std=gnu++98") + set(CMAKE_HIP98_STANDARD__HAS_FULL_SUPPORT ON) + + set(CMAKE_HIP11_STANDARD_COMPILE_OPTION "-std=c++11") + set(CMAKE_HIP11_EXTENSION_COMPILE_OPTION "-std=gnu++11") + set(CMAKE_HIP11_STANDARD__HAS_FULL_SUPPORT ON) + + set(CMAKE_HIP14_STANDARD_COMPILE_OPTION "-std=c++14") + set(CMAKE_HIP14_EXTENSION_COMPILE_OPTION "-std=gnu++14") + set(CMAKE_HIP14_STANDARD__HAS_FULL_SUPPORT ON) + + set(CMAKE_HIP17_STANDARD_COMPILE_OPTION "-std=c++17") + set(CMAKE_HIP17_EXTENSION_COMPILE_OPTION "-std=gnu++17") + set(CMAKE_HIP17_STANDARD__HAS_FULL_SUPPORT ON) + + set(CMAKE_HIP20_STANDARD_COMPILE_OPTION "-std=c++20") + set(CMAKE_HIP20_EXTENSION_COMPILE_OPTION "-std=gnu++20") +endif() + +set(CMAKE_HIP_RUNTIME_LIBRARY_DEFAULT "SHARED") +set(CMAKE_HIP_RUNTIME_LIBRARY_LINK_OPTIONS_STATIC "") +set(CMAKE_HIP_RUNTIME_LIBRARY_LINK_OPTIONS_SHARED "") + +__compiler_check_default_language_standard(HIP 3.5 11) 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/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 const& argv, LanguageStandardState cState("C"); LanguageStandardState cudaState("CUDA"); LanguageStandardState cxxState("CXX"); + LanguageStandardState hipState("HIP"); LanguageStandardState objcState("OBJC"); LanguageStandardState objcxxState("OBJCXX"); std::vector targets; @@ -323,6 +328,7 @@ int cmCoreTryCompile::TryCompileCode(std::vector 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 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 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 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 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 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 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 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 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 6209f3e..812535b 100644 --- a/Source/cmGeneratorTarget.cxx +++ b/Source/cmGeneratorTarget.cxx @@ -980,7 +980,7 @@ cmProp cmGeneratorTarget::GetPropertyWithPairedLanguageSupport( // Check if we should use the value set by another language. if (lang == "OBJC") { propertyValue = this->GetPropertyWithPairedLanguageSupport("C", suffix); - } else if (lang == "OBJCXX" || lang == "CUDA") { + } else if (lang == "OBJCXX" || lang == "CUDA" || lang == "HIP") { propertyValue = this->GetPropertyWithPairedLanguageSupport("CXX", suffix); } @@ -1121,7 +1121,7 @@ void cmGeneratorTarget::AppendLanguageSideEffects( std::map>& sideEffects) const { static const std::set LANGS_WITH_NO_SIDE_EFFECTS = { - "C"_s, "CXX"_s, "OBJC"_s, "OBJCXX"_s, "ASM"_s, "CUDA"_s, + "C"_s, "CXX"_s, "OBJC"_s, "OBJCXX"_s, "ASM"_s, "CUDA"_s, "HIP"_s }; for (auto const& lang : this->GetAllConfigCompileLanguages()) { @@ -3351,6 +3351,23 @@ void cmGeneratorTarget::AddISPCTargetFlags(std::string& flags) const } } +void cmGeneratorTarget::AddHIPArchitectureFlags(std::string& flags) const +{ + const std::string& property = this->GetSafeProperty("HIP_ARCHITECTURES"); + + // If HIP_ARCHITECTURES is false we don't add any architectures. + if (cmIsOff(property)) { + return; + } + + std::vector options; + cmExpandList(property, options); + + for (std::string& option : options) { + flags += " --offload-arch=" + option; + } +} + void cmGeneratorTarget::AddCUDAToolkitFlags(std::string& flags) const { std::string const& compiler = @@ -4742,7 +4759,8 @@ bool cmGeneratorTarget::ComputeCompileFeatures( } // Custom updates for the CUDA standard. - if (generatorTargetLanguageStandard && language.first == "CUDA") { + if (generatorTargetLanguageStandard != nullptr && + (language.first == "CUDA")) { if (generatorTargetLanguageStandard->Value == "98") { this->LanguageStandardMap[key].Value = "03"; } diff --git a/Source/cmGeneratorTarget.h b/Source/cmGeneratorTarget.h index be36f3f..ab9ccfd 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( 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::vector pairedLanguages{ { "OBJC", "C" }, - { "OBJCXX", "CXX" }, - { "CUDA", "CXX" } }; + std::vector pairedLanguages{ + { "OBJC", "C" }, { "OBJCXX", "CXX" }, { "CUDA", "CXX" }, { "HIP", "CXX" } + }; std::set 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 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(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 StandardComputerMapping = - { - { "C", + { { "C", StanardLevelComputer{ "C", std::vector{ 90, 99, 11, 17, 23 }, std::vector{ "90", "99", "11", "17", "23" } } }, @@ -326,7 +328,10 @@ std::unordered_map StandardComputerMapping = StanardLevelComputer{ "OBJCXX", std::vector{ 98, 11, 14, 17, 20, 23 }, std::vector{ "98", "11", "14", "17", "20", "23" } } }, - }; + { "HIP", + StanardLevelComputer{ + "HIP", std::vector{ 98, 11, 14, 17, 20, 23 }, + std::vector{ "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(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 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 38b4301..346c2c5 100644 --- a/Tests/CMakeLists.txt +++ b/Tests/CMakeLists.txt @@ -1479,6 +1479,10 @@ if(BUILD_TESTING) add_subdirectory(CudaOnly) endif() + if(CMake_TEST_HIP) + add_subdirectory(HIP) + endif() + if(CMake_TEST_ISPC) add_subdirectory(ISPC) endif() diff --git a/Tests/HIP/ArchitectureOff/CMakeLists.txt b/Tests/HIP/ArchitectureOff/CMakeLists.txt new file mode 100644 index 0000000..bccb3b4 --- /dev/null +++ b/Tests/HIP/ArchitectureOff/CMakeLists.txt @@ -0,0 +1,8 @@ +cmake_minimum_required(VERSION 3.18) +project(HIPArchitecture HIP) + +# Make sure CMake doesn't pass architectures if HIP_ARCHITECTURES is OFF. +string(APPEND CMAKE_HIP_FLAGS " --offload-arch=gfx908") + +add_executable(HIPOnlyArchitectureOff main.hip) +set_property(TARGET HIPOnlyArchitectureOff PROPERTY HIP_ARCHITECTURES OFF) diff --git a/Tests/HIP/ArchitectureOff/main.hip b/Tests/HIP/ArchitectureOff/main.hip new file mode 100644 index 0000000..9256318 --- /dev/null +++ b/Tests/HIP/ArchitectureOff/main.hip @@ -0,0 +1,9 @@ +#ifdef __HIP_DEVICE_COMPILE__ +# ifndef __gfx908__ +# error "Passed architecture gfx908, but got something else." +# endif +#endif + +int main() +{ +} diff --git a/Tests/HIP/CMakeLists.txt b/Tests/HIP/CMakeLists.txt new file mode 100644 index 0000000..b3966c9 --- /dev/null +++ b/Tests/HIP/CMakeLists.txt @@ -0,0 +1,12 @@ +macro (add_hip_test_macro name) + add_test_macro("${name}" ${ARGN}) + set_property(TEST "${name}" APPEND + PROPERTY LABELS "HIP") +endmacro () + +add_hip_test_macro(HIP.ArchitectureOff HIPOnlyArchitectureOff) +add_hip_test_macro(HIP.CompileFlags HIPOnlyCompileFlags) +add_hip_test_macro(HIP.EnableStandard HIPEnableStandard) +add_hip_test_macro(HIP.MathFunctions HIPOnlyMathFunctions) +add_hip_test_macro(HIP.TryCompile HIPOnlyTryCompile) +add_hip_test_macro(HIP.WithDefs HIPOnlyWithDefs) diff --git a/Tests/HIP/CompileFlags/CMakeLists.txt b/Tests/HIP/CompileFlags/CMakeLists.txt new file mode 100644 index 0000000..c808313 --- /dev/null +++ b/Tests/HIP/CompileFlags/CMakeLists.txt @@ -0,0 +1,8 @@ +cmake_minimum_required(VERSION 3.18) +project(CompileFlags HIP) + +add_executable(HIPOnlyCompileFlags main.hip) + +set_property(TARGET HIPOnlyCompileFlags PROPERTY HIP_ARCHITECTURES gfx803) + +target_compile_options(HIPOnlyCompileFlags PRIVATE -DALWAYS_DEFINE) diff --git a/Tests/HIP/CompileFlags/main.hip b/Tests/HIP/CompileFlags/main.hip new file mode 100644 index 0000000..3259f1d --- /dev/null +++ b/Tests/HIP/CompileFlags/main.hip @@ -0,0 +1,13 @@ +#ifdef __HIP_DEVICE_COMPILE__ +# ifndef __gfx803__ +# error "Passed architecture gfx803, but got something else." +# endif +#endif + +#ifndef ALWAYS_DEFINE +# error "ALWAYS_DEFINE not defined!" +#endif + +int main() +{ +} diff --git a/Tests/HIP/EnableStandard/CMakeLists.txt b/Tests/HIP/EnableStandard/CMakeLists.txt new file mode 100644 index 0000000..6701724 --- /dev/null +++ b/Tests/HIP/EnableStandard/CMakeLists.txt @@ -0,0 +1,20 @@ +cmake_minimum_required(VERSION 3.18) +project (EnableStandard HIP) + +set(CMAKE_CXX_COMPILER ${CMAKE_HIP_COMPILER}) +enable_language(CXX) + +#Goal for this example: +#build hip sources that require C++11 to be enabled. + +add_library(HIPStatic11 STATIC static.cxx) +set_source_files_properties(static.cxx PROPERTIES LANGUAGE HIP) + +add_library(HIPDynamic11 SHARED shared.hip) + +add_executable(HIPEnableStandard main.hip) +target_link_libraries(HIPEnableStandard PRIVATE HIPStatic11 HIPDynamic11) + +target_compile_features(HIPDynamic11 PRIVATE cxx_std_11) +set_target_properties(HIPStatic11 PROPERTIES HIP_STANDARD 11) +set_target_properties(HIPStatic11 PROPERTIES HIP_STANDARD_REQUIRED TRUE) diff --git a/Tests/HIP/EnableStandard/main.hip b/Tests/HIP/EnableStandard/main.hip new file mode 100644 index 0000000..7dc32e8 --- /dev/null +++ b/Tests/HIP/EnableStandard/main.hip @@ -0,0 +1,23 @@ + +#include + +#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 + +#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::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 + +using tt = std::true_type; +using ft = std::false_type; +int __host__ static_hip11_func(int x) +{ + return x * x + std::integral_constant::value; +} diff --git a/Tests/HIP/MathFunctions/CMakeLists.txt b/Tests/HIP/MathFunctions/CMakeLists.txt new file mode 100644 index 0000000..81e3ddb --- /dev/null +++ b/Tests/HIP/MathFunctions/CMakeLists.txt @@ -0,0 +1,18 @@ +cmake_minimum_required(VERSION 3.18) +project(MathFunctions HIP) + +# This test covers these major HIP language/runtime requirements: +# +# 1. This makes sure CMake properly specifies the internal clang header dirs +# that hold headers needed for overloads of device side functions +# +# 2. This makes sure that all HIP include directories are properly marked as +# system includes so we don't get the following warnings: +# replacement function 'operator delete' cannot be declared 'inline'# +# +# 3. This makes sure CMake properly links to all the built-in libraries +# that hip needs that inject support for __half support +# +add_executable(HIPOnlyMathFunctions main.hip) +target_compile_options(HIPOnlyMathFunctions PRIVATE -Werror) +target_compile_features(HIPOnlyMathFunctions PRIVATE hip_std_14) diff --git a/Tests/HIP/MathFunctions/main.hip b/Tests/HIP/MathFunctions/main.hip new file mode 100644 index 0000000..8a6e77f --- /dev/null +++ b/Tests/HIP/MathFunctions/main.hip @@ -0,0 +1,40 @@ + +#include +#include +#include +#include + +#include +#include + +namespace { +template +__global__ void global_entry_point(F f, T *out) { + *out = f(); +} + +template +bool verify(F f, T expected) +{ + std::unique_ptr cpu_T(new T); + T* gpu_T = nullptr; + hipMalloc((void**)&gpu_T, sizeof(T)); + hipLaunchKernelGGL(global_entry_point, 1, 1, 0, 0, f, gpu_T); + hipMemcpy(cpu_T.get(), gpu_T, sizeof(T), hipMemcpyDeviceToHost); + hipFree(gpu_T); + return (*cpu_T == expected); +} +} + +int main(int argc, char** argv) +{ + bool valid = verify([]__device__(){ return std::round(1.4f); }, 1.0f); + valid &= verify([]__device__(){ return max<_Float16>(1.0f, 2.0f); }, 2.0f); + valid &= verify([]__device__(){ return min<_Float16>(1.0f, 2.0f); }, 1.0f); + + if (valid) { + return 0; + } else { + return 1; + } +} diff --git a/Tests/HIP/TryCompile/CMakeLists.txt b/Tests/HIP/TryCompile/CMakeLists.txt new file mode 100644 index 0000000..f3bb3bf --- /dev/null +++ b/Tests/HIP/TryCompile/CMakeLists.txt @@ -0,0 +1,14 @@ +cmake_minimum_required(VERSION 3.18) +project (TryCompile HIP) + +#Goal for this example: +# Verify try_compile with HIP language works +set(CMAKE_HIP_STANDARD 14) + +set(CMAKE_TRY_COMPILE_TARGET_TYPE STATIC_LIBRARY) +try_compile(result "${CMAKE_CURRENT_BINARY_DIR}" + SOURCES "${CMAKE_CURRENT_SOURCE_DIR}/device_function.hip" + COPY_FILE "${CMAKE_CURRENT_BINARY_DIR}/device_function.o") + +add_executable(HIPOnlyTryCompile main.hip) +target_link_libraries(HIPOnlyTryCompile "${CMAKE_CURRENT_BINARY_DIR}/device_function.o") diff --git a/Tests/HIP/TryCompile/device_function.hip b/Tests/HIP/TryCompile/device_function.hip new file mode 100644 index 0000000..7c1205e --- /dev/null +++ b/Tests/HIP/TryCompile/device_function.hip @@ -0,0 +1,17 @@ +#include +#include + +static __global__ void fake_hip_kernel() +{ +} + +int __host__ try_compile_hip_func(int x) +{ + + fake_hip_kernel<<<1, 1>>>(); + bool valid = (hipSuccess == hipGetLastError()); + if (!valid) { + throw std::system_error(ENODEV, std::generic_category(), "no hip device"); + } + return x * x; +} diff --git a/Tests/HIP/TryCompile/main.hip b/Tests/HIP/TryCompile/main.hip new file mode 100644 index 0000000..091dca3 --- /dev/null +++ b/Tests/HIP/TryCompile/main.hip @@ -0,0 +1,8 @@ +int __host__ try_compile_hip_func(int x); + +int main(int argc, char** argv) +{ + try_compile_hip_func(int(42)); + + return 0; +} diff --git a/Tests/HIP/WithDefs/CMakeLists.txt b/Tests/HIP/WithDefs/CMakeLists.txt new file mode 100644 index 0000000..3d4461b --- /dev/null +++ b/Tests/HIP/WithDefs/CMakeLists.txt @@ -0,0 +1,36 @@ + +cmake_minimum_required(VERSION 3.18) +project (WithDefs HIP) + +set(CMAKE_HIP_ARCHITECTURES OFF) +set(release_compile_defs DEFREL) + +#Goal for this example: +#build a executable that needs to be passed a complex define through add_definitions +#this verifies we can pass C++ style attributes to hipcc +add_definitions("-DPACKED_DEFINE=[[gnu::packed]]") + +add_executable(HIPOnlyWithDefs main.hip.cpp) +set_source_files_properties(main.hip.cpp PROPERTIES LANGUAGE HIP) + +target_compile_features(HIPOnlyWithDefs PRIVATE hip_std_17) + +target_compile_options(HIPOnlyWithDefs + PRIVATE + -DFLAG_COMPILE_LANG_$ + $<$:-DFLAG_LANG_IS_HIP=$> # Host-only defines are possible only on NVCC. + ) + +target_compile_definitions(HIPOnlyWithDefs + PRIVATE + $<$:$> + -DDEF_COMPILE_LANG_$ + -DDEF_LANG_IS_HIP=$ + -DDEF_HIP_COMPILER=$ + -DDEF_HIP_COMPILER_VERSION=$ + ) + +target_include_directories(HIPOnlyWithDefs + PRIVATE + $<$:${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 + +#include +#include +#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 d02d7a2..16487f3 100644 --- a/Tests/RunCMake/CMakeLists.txt +++ b/Tests/RunCMake/CMakeLists.txt @@ -724,6 +724,9 @@ if("${CMAKE_GENERATOR}" MATCHES "Make|Ninja") if(DEFINED CMake_TEST_CUDA) list(APPEND CompilerLauncher_ARGS -DCMake_TEST_CUDA=${CMake_TEST_CUDA}) endif() + if(DEFINED CMake_TEST_HIP) + list(APPEND CompilerLauncher_ARGS -DCMake_TEST_HIP=${CMake_TEST_HIP}) + endif() if(DEFINED CMake_TEST_ISPC) list(APPEND CompilerLauncher_ARGS -DCMake_TEST_ISPC=${CMake_TEST_ISPC}) endif() @@ -736,7 +739,7 @@ if("${CMAKE_GENERATOR}" MATCHES "Make|Ninja") endif() add_RunCMake_test(CompilerLauncher) set_property(TEST RunCMake.CompilerLauncher APPEND - PROPERTY LABELS "CUDA;ISPC") + PROPERTY LABELS "CUDA;HIP;ISPC") add_RunCMake_test(ctest_labels_for_subprojects) add_RunCMake_test(CompilerArgs) add_RunCMake_test(LinkerLauncher) 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; +} -- cgit v0.12