From ed3aad97e3d3cb27637722fabb0ae3a5611daf5c Mon Sep 17 00:00:00 2001 From: Brad King Date: Fri, 4 Jun 2021 12:13:29 -0400 Subject: cmGeneratorTarget: Clear cached link implementation when a source is added Adding a source file at generate time can cause the linker language or other settings to change that affect `GetLinkImplementationLibraries` and friends. --- Source/cmGeneratorTarget.cxx | 1 + 1 file changed, 1 insertion(+) diff --git a/Source/cmGeneratorTarget.cxx b/Source/cmGeneratorTarget.cxx index 5deb2df..6209f3e 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) -- cgit v0.12 From 590553f32246b374853c50cf7c30af3880ac02ea Mon Sep 17 00:00:00 2001 From: Zack Galbreath Date: Mon, 7 Jun 2021 18:18:29 +0000 Subject: Compilers: protect use of __has_include If the compiler does not have __has_include, pretend the answer is always no --- Modules/CMakeCCompilerId.c.in | 6 ++++++ Modules/CMakeCXXCompilerId.cpp.in | 6 ++++++ 2 files changed, 12 insertions(+) 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 -- cgit v0.12 From bd844387df133897ce6e3f290af771acb0df4df9 Mon Sep 17 00:00:00 2001 From: Robert Maynard Date: Fri, 28 Aug 2020 15:00:57 -0400 Subject: ROCMClang: Add the ROCm toolkit derived clang compiler to CMake --- Help/variable/CMAKE_LANG_COMPILER_ID.rst | 1 + Modules/CMakeCompilerIdDetection.cmake | 5 ++++ Modules/CMakeDetermineCompilerId.cmake | 1 - Modules/Compiler/ROCMClang-ASM.cmake | 2 ++ Modules/Compiler/ROCMClang-C.cmake | 7 +++++ Modules/Compiler/ROCMClang-CXX.cmake | 7 +++++ Modules/Compiler/ROCMClang-DetermineCompiler.cmake | 19 ++++++++++++ Modules/Compiler/ROCMClang-FindBinUtils.cmake | 1 + Modules/Compiler/ROCMClang-OBJC.cmake | 7 +++++ Modules/Compiler/ROCMClang-OBJCXX.cmake | 7 +++++ Modules/Compiler/ROCMClang.cmake | 35 ++++++++++++++++++++++ Modules/WriteCompilerDetectionHeader.cmake | 1 + 12 files changed, 92 insertions(+), 1 deletion(-) create mode 100644 Modules/Compiler/ROCMClang-ASM.cmake create mode 100644 Modules/Compiler/ROCMClang-C.cmake create mode 100644 Modules/Compiler/ROCMClang-CXX.cmake create mode 100644 Modules/Compiler/ROCMClang-DetermineCompiler.cmake create mode 100644 Modules/Compiler/ROCMClang-FindBinUtils.cmake create mode 100644 Modules/Compiler/ROCMClang-OBJC.cmake create mode 100644 Modules/Compiler/ROCMClang-OBJCXX.cmake create mode 100644 Modules/Compiler/ROCMClang.cmake 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/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/CMakeDetermineCompilerId.cmake b/Modules/CMakeDetermineCompilerId.cmake index 448f071..d656cd0 100644 --- a/Modules/CMakeDetermineCompilerId.cmake +++ b/Modules/CMakeDetermineCompilerId.cmake @@ -150,7 +150,6 @@ function(CMAKE_DETERMINE_COMPILER_ID lang flagvar src) endif() endif() - if (COMPILER_QNXNTO AND CMAKE_${lang}_COMPILER_ID STREQUAL "GNU") execute_process( COMMAND "${CMAKE_${lang}_COMPILER}" 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()") + +set(_compiler_id_version_compute " +# if defined(__clang__) && __has_include() +# include +# 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-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/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) -- cgit v0.12 From ff0d2858e1e47af8e849318b411b1fbd8579a053 Mon Sep 17 00:00:00 2001 From: Robert Maynard Date: Wed, 21 Oct 2020 14:24:37 -0400 Subject: HIP: Extract clang compiler details from hipcc --- Modules/CMakeDetermineCompilerId.cmake | 35 ++++++++++++++++++++++++++++++++++ 1 file changed, 35 insertions(+) diff --git a/Modules/CMakeDetermineCompilerId.cmake b/Modules/CMakeDetermineCompilerId.cmake index d656cd0..bd1e732 100644 --- a/Modules/CMakeDetermineCompilerId.cmake +++ b/Modules/CMakeDetermineCompilerId.cmake @@ -150,6 +150,41 @@ 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( COMMAND "${CMAKE_${lang}_COMPILER}" -- cgit v0.12 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 From 947dbed0aa502c70a9e165757450bd8409036980 Mon Sep 17 00:00:00 2001 From: Robert Maynard Date: Wed, 2 Sep 2020 14:53:41 -0400 Subject: HIP: Automatically inject the `hip::device` runtime target Any target that might need to link to hip code needs the `hip::device` target --- Modules/CMakeHIPCompiler.cmake.in | 5 +- Modules/CMakeHIPInformation.cmake | 7 ++ Modules/CMakeHIPRuntime.cmake.in | 99 +++++++++++++++++ Modules/CMakeTestHIPCompiler.cmake | 21 ++++ Modules/Compiler/Clang-HIP.cmake | 4 + Modules/Compiler/ROCMClang-HIP.cmake | 4 + Source/cmComputeLinkDepends.cxx | 14 +++ Source/cmGeneratorTarget.cxx | 196 ++++++++++++++++++++++++++++----- Source/cmGeneratorTarget.h | 6 + Source/cmLinkItem.h | 5 + Tests/HIP/CMakeLists.txt | 3 + Tests/HIP/InferHipLang1/CMakeLists.txt | 12 ++ Tests/HIP/InferHipLang1/interface.hip | 19 ++++ Tests/HIP/InferHipLang1/main.cxx | 19 ++++ Tests/HIP/InferHipLang2/CMakeLists.txt | 12 ++ Tests/HIP/InferHipLang2/interface.hip | 19 ++++ Tests/HIP/InferHipLang2/main.cxx | 19 ++++ Tests/HIP/MixedLanguage/CMakeLists.txt | 19 ++++ Tests/HIP/MixedLanguage/main.cxx | 40 +++++++ Tests/HIP/MixedLanguage/shared.c | 12 ++ Tests/HIP/MixedLanguage/shared.cxx | 21 ++++ Tests/HIP/MixedLanguage/shared.hip | 26 +++++ Tests/HIP/MixedLanguage/static.c | 6 + Tests/HIP/MixedLanguage/static.cxx | 7 ++ Tests/HIP/MixedLanguage/static.hip | 21 ++++ 25 files changed, 587 insertions(+), 29 deletions(-) create mode 100644 Modules/CMakeHIPRuntime.cmake.in create mode 100644 Tests/HIP/InferHipLang1/CMakeLists.txt create mode 100644 Tests/HIP/InferHipLang1/interface.hip create mode 100644 Tests/HIP/InferHipLang1/main.cxx create mode 100644 Tests/HIP/InferHipLang2/CMakeLists.txt create mode 100644 Tests/HIP/InferHipLang2/interface.hip create mode 100644 Tests/HIP/InferHipLang2/main.cxx create mode 100644 Tests/HIP/MixedLanguage/CMakeLists.txt create mode 100644 Tests/HIP/MixedLanguage/main.cxx create mode 100644 Tests/HIP/MixedLanguage/shared.c create mode 100644 Tests/HIP/MixedLanguage/shared.cxx create mode 100644 Tests/HIP/MixedLanguage/shared.hip create mode 100644 Tests/HIP/MixedLanguage/static.c create mode 100644 Tests/HIP/MixedLanguage/static.cxx create mode 100644 Tests/HIP/MixedLanguage/static.hip diff --git a/Modules/CMakeHIPCompiler.cmake.in b/Modules/CMakeHIPCompiler.cmake.in index eeb51b2..9a30a45 100644 --- a/Modules/CMakeHIPCompiler.cmake.in +++ b/Modules/CMakeHIPCompiler.cmake.in @@ -22,7 +22,7 @@ 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 90) set(CMAKE_HIP_LINKER_PREFERENCE_PROPAGATES 1) set(CMAKE_HIP_SIZEOF_DATA_PTR "@CMAKE_HIP_SIZEOF_DATA_PTR@") @@ -48,8 +48,7 @@ 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_HIP_RUNTIME_LIBRARY_DEFAULT "SHARED") set(CMAKE_AR "@CMAKE_AR@") set(CMAKE_HIP_COMPILER_AR "@CMAKE_HIP_COMPILER_AR@") diff --git a/Modules/CMakeHIPInformation.cmake b/Modules/CMakeHIPInformation.cmake index df1d98a..ec37e1c 100644 --- a/Modules/CMakeHIPInformation.cmake +++ b/Modules/CMakeHIPInformation.cmake @@ -130,3 +130,10 @@ if(NOT CMAKE_HIP_LINK_EXECUTABLE) 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 index da70f8d..d9fcc9d 100644 --- a/Modules/CMakeTestHIPCompiler.cmake +++ b/Modules/CMakeTestHIPCompiler.cmake @@ -78,6 +78,20 @@ unset(__CMAKE_HIP_FLAGS) 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 @@ -96,3 +110,10 @@ 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/Compiler/Clang-HIP.cmake b/Modules/Compiler/Clang-HIP.cmake index 6ac39de..1030a43 100644 --- a/Modules/Compiler/Clang-HIP.cmake +++ b/Modules/Compiler/Clang-HIP.cmake @@ -14,3 +14,7 @@ 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-HIP.cmake b/Modules/Compiler/ROCMClang-HIP.cmake index 9e792b5..7af7699 100644 --- a/Modules/Compiler/ROCMClang-HIP.cmake +++ b/Modules/Compiler/ROCMClang-HIP.cmake @@ -42,4 +42,8 @@ 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/Source/cmComputeLinkDepends.cxx b/Source/cmComputeLinkDepends.cxx index 4a6518fd..0b27e34 100644 --- a/Source/cmComputeLinkDepends.cxx +++ b/Source/cmComputeLinkDepends.cxx @@ -7,6 +7,7 @@ #include #include #include +#include #include #include @@ -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/cmGeneratorTarget.cxx b/Source/cmGeneratorTarget.cxx index 812535b..cb10117 100644 --- a/Source/cmGeneratorTarget.cxx +++ b/Source/cmGeneratorTarget.cxx @@ -1235,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()); @@ -1474,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 const& libraries) +{ + for (cmLinkImplItem const& lib : libraries) { + if (lib.Target) { + EvaluatedTargetPropertyEntry ee(lib, lib.Backtrace); + // Pretend $ 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 $ 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); } } } @@ -1656,7 +1719,8 @@ std::vector> 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::size_type numFilesBefore = files.size(); bool contextDependentInterfaceSources = processSources( this, linkInterfaceSourcesEntries, files, uniqueSrcs, debugSources); @@ -3587,7 +3651,7 @@ std::vector> cmGeneratorTarget::GetIncludeDirectories( } AddInterfaceEntries(this, config, "INTERFACE_INCLUDE_DIRECTORIES", lang, - &dagChecker, entries); + &dagChecker, entries, IncludeRuntimeInterface::Yes); if (this->Makefile->IsOn("APPLE")) { if (cmLinkImplementationLibraries const* impl = @@ -3812,7 +3876,7 @@ std::vector> 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); @@ -3854,7 +3918,8 @@ std::vector> 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); @@ -3898,7 +3963,7 @@ std::vector> 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 = @@ -3956,7 +4021,7 @@ std::vector> cmGeneratorTarget::GetPrecompileHeaders( this, config, language, &dagChecker, this->PrecompileHeadersEntries); AddInterfaceEntries(this, config, "INTERFACE_PRECOMPILE_HEADERS", language, - &dagChecker, entries); + &dagChecker, entries, IncludeRuntimeInterface::Yes); std::vector> list; processOptions(this, entries, list, uniqueOptions, debugDefines, @@ -4341,7 +4406,7 @@ std::vector> 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, @@ -4600,7 +4665,7 @@ std::vector> 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, @@ -4639,7 +4704,7 @@ std::vector> 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", @@ -6382,6 +6447,7 @@ cmLinkInterface const* cmGeneratorTarget::GetLinkInterface( iface.AllDone = true; if (iface.Exists) { this->ComputeLinkInterface(config, iface, head, secondPass); + this->ComputeLinkInterfaceRuntimeLibraries(config, iface); } } @@ -6905,6 +6971,83 @@ void cmGeneratorTarget::ComputeLinkInterfaceLibraries( } } +namespace { + +template +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 +std::vector computeImplicitLanguageTargets( + std::string const& lang, std::string const& config, + cmGeneratorTarget const* currentTarget) +{ + cmListFileBacktrace bt; + std::vector 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 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(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(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(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 @@ -7169,6 +7312,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 ab9ccfd..ed66fb1 100644 --- a/Source/cmGeneratorTarget.h +++ b/Source/cmGeneratorTarget.h @@ -1111,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& 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 #include #include +#include #include #include @@ -80,6 +81,8 @@ struct cmLinkInterface : public cmLinkInterfaceLibraries { // Languages whose runtime libraries must be linked. std::vector Languages; + std::unordered_map> + LanguageRuntimeLibraries; // Shared library dependencies needed for linking on some platforms. std::vector SharedDeps; @@ -115,6 +118,8 @@ struct cmLinkImplementation : public cmLinkImplementationLibraries { // Languages whose runtime libraries must be linked. std::vector Languages; + std::unordered_map> + LanguageRuntimeLibraries; // Whether the list depends on a link language genex. bool HadLinkLanguageSensitiveCondition = false; diff --git a/Tests/HIP/CMakeLists.txt b/Tests/HIP/CMakeLists.txt index b3966c9..9499be8 100644 --- a/Tests/HIP/CMakeLists.txt +++ b/Tests/HIP/CMakeLists.txt @@ -7,6 +7,9 @@ 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/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 +#include +#include + +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::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 + +#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 +#include +#include + +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::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 + +#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/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 + +#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 + +#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::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 +#include +#include + +#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::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 + +int static_cxx_func(int x) +{ + return x * x + std::integral_constant::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 +#include +#include + + +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::value; +} -- cgit v0.12 From 2e86e50c2fb2897dc1a0612676c5bd26eb7924a9 Mon Sep 17 00:00:00 2001 From: Robert Maynard Date: Sun, 25 Oct 2020 14:34:18 -0400 Subject: HIP: Add HIP to all the Check* modules --- Modules/CheckLinkerFlag.cmake | 2 ++ Modules/Internal/CheckCompilerFlag.cmake | 3 +++ Modules/Internal/CheckSourceCompiles.cmake | 3 +++ Modules/Internal/CheckSourceRuns.cmake | 3 +++ Tests/RunCMake/CMakeLists.txt | 18 +++++++++++---- .../CheckCompilerFlag/CheckHIPCompilerFlag.cmake | 13 +++++++++++ .../RunCMake/CheckCompilerFlag/RunCMakeTest.cmake | 4 ++++ .../CheckLinkerFlag/CheckHIPLinkerFlag.cmake | 3 +++ Tests/RunCMake/CheckLinkerFlag/RunCMakeTest.cmake | 4 ++++ .../CheckHIPSourceCompiles.cmake | 27 ++++++++++++++++++++++ .../CheckSourceCompiles/RunCMakeTest.cmake | 4 ++++ .../CheckSourceRuns/CheckHIPSourceRuns.cmake | 21 +++++++++++++++++ Tests/RunCMake/CheckSourceRuns/RunCMakeTest.cmake | 4 ++++ 13 files changed, 105 insertions(+), 4 deletions(-) create mode 100644 Tests/RunCMake/CheckCompilerFlag/CheckHIPCompilerFlag.cmake create mode 100644 Tests/RunCMake/CheckLinkerFlag/CheckHIPLinkerFlag.cmake create mode 100644 Tests/RunCMake/CheckSourceCompiles/CheckHIPSourceCompiles.cmake create mode 100644 Tests/RunCMake/CheckSourceRuns/CheckHIPSourceRuns.cmake 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/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/Tests/RunCMake/CMakeLists.txt b/Tests/RunCMake/CMakeLists.txt index 16487f3..dcf1e72 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 @@ -600,12 +601,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 @@ -613,14 +617,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() 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 + __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 + __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() -- cgit v0.12 From 20d086f1a2c320da59830682a19e006ec19dca94 Mon Sep 17 00:00:00 2001 From: Robert Maynard Date: Mon, 26 Oct 2020 21:10:18 +0000 Subject: HIP: All HIP tests now run on CMake's current AMD hardware --- Tests/HIP/InferHipLang1/CMakeLists.txt | 1 + Tests/HIP/InferHipLang2/CMakeLists.txt | 1 + Tests/HIP/MathFunctions/CMakeLists.txt | 1 + Tests/HIP/MixedLanguage/CMakeLists.txt | 1 + Tests/HIP/TryCompile/CMakeLists.txt | 1 + Tests/HIP/WithDefs/CMakeLists.txt | 1 + 6 files changed, 6 insertions(+) diff --git a/Tests/HIP/InferHipLang1/CMakeLists.txt b/Tests/HIP/InferHipLang1/CMakeLists.txt index 63d77fd..f3c5421 100644 --- a/Tests/HIP/InferHipLang1/CMakeLists.txt +++ b/Tests/HIP/InferHipLang1/CMakeLists.txt @@ -10,3 +10,4 @@ target_compile_features(InterfaceWithHIP INTERFACE cxx_std_11) add_executable(HIPInferHipLang1 ) target_link_libraries(HIPInferHipLang1 PRIVATE InterfaceWithHIP) +set_property(TARGET HIPInferHipLang1 PROPERTY HIP_ARCHITECTURES gfx900) diff --git a/Tests/HIP/InferHipLang2/CMakeLists.txt b/Tests/HIP/InferHipLang2/CMakeLists.txt index 0e69de3..a8a55b7 100644 --- a/Tests/HIP/InferHipLang2/CMakeLists.txt +++ b/Tests/HIP/InferHipLang2/CMakeLists.txt @@ -7,6 +7,7 @@ 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) +set_property(TARGET InterfaceWithHIP PROPERTY HIP_ARCHITECTURES gfx900) add_executable(HIPInferHipLang2 ) target_link_libraries(HIPInferHipLang2 PRIVATE InterfaceWithHIP) diff --git a/Tests/HIP/MathFunctions/CMakeLists.txt b/Tests/HIP/MathFunctions/CMakeLists.txt index 81e3ddb..69d63dd 100644 --- a/Tests/HIP/MathFunctions/CMakeLists.txt +++ b/Tests/HIP/MathFunctions/CMakeLists.txt @@ -13,6 +13,7 @@ project(MathFunctions HIP) # 3. This makes sure CMake properly links to all the built-in libraries # that hip needs that inject support for __half support # +set(CMAKE_HIP_ARCHITECTURES "gfx900") add_executable(HIPOnlyMathFunctions main.hip) target_compile_options(HIPOnlyMathFunctions PRIVATE -Werror) target_compile_features(HIPOnlyMathFunctions PRIVATE hip_std_14) diff --git a/Tests/HIP/MixedLanguage/CMakeLists.txt b/Tests/HIP/MixedLanguage/CMakeLists.txt index 4f6dd3b..ce2ff89 100644 --- a/Tests/HIP/MixedLanguage/CMakeLists.txt +++ b/Tests/HIP/MixedLanguage/CMakeLists.txt @@ -3,6 +3,7 @@ project (MixedLanguage C CXX HIP) set(CMAKE_HIP_STANDARD 14) set(CMAKE_CXX_STANDARD 14) +set(CMAKE_HIP_ARCHITECTURES "gfx900") #Goal for this example: #make sure that we can build multiple languages into targets diff --git a/Tests/HIP/TryCompile/CMakeLists.txt b/Tests/HIP/TryCompile/CMakeLists.txt index f3bb3bf..92a834c 100644 --- a/Tests/HIP/TryCompile/CMakeLists.txt +++ b/Tests/HIP/TryCompile/CMakeLists.txt @@ -4,6 +4,7 @@ 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}" diff --git a/Tests/HIP/WithDefs/CMakeLists.txt b/Tests/HIP/WithDefs/CMakeLists.txt index 3d4461b..e2db182 100644 --- a/Tests/HIP/WithDefs/CMakeLists.txt +++ b/Tests/HIP/WithDefs/CMakeLists.txt @@ -17,6 +17,7 @@ target_compile_features(HIPOnlyWithDefs PRIVATE hip_std_17) target_compile_options(HIPOnlyWithDefs PRIVATE + --offload-arch=gfx900 -DFLAG_COMPILE_LANG_$ $<$:-DFLAG_LANG_IS_HIP=$> # Host-only defines are possible only on NVCC. ) -- cgit v0.12 From 8514ee9b315b4ad02eed0e3cf11d8579f307fac0 Mon Sep 17 00:00:00 2001 From: Zack Galbreath Date: Thu, 27 May 2021 18:03:51 +0000 Subject: HIP: analyze output of `hipcc` to determine default GPU architecture --- Help/prop_tgt/HIP_ARCHITECTURES.rst | 5 +++-- Help/variable/CMAKE_HIP_ARCHITECTURES.rst | 2 +- Modules/CMakeDetermineHIPCompiler.cmake | 7 ++++++- Source/cmGeneratorTarget.cxx | 6 ++++++ Tests/HIP/InferHipLang1/CMakeLists.txt | 1 - Tests/HIP/InferHipLang2/CMakeLists.txt | 1 - Tests/HIP/MathFunctions/CMakeLists.txt | 1 - Tests/HIP/MixedLanguage/CMakeLists.txt | 1 - 8 files changed, 16 insertions(+), 8 deletions(-) diff --git a/Help/prop_tgt/HIP_ARCHITECTURES.rst b/Help/prop_tgt/HIP_ARCHITECTURES.rst index 052f472..06f956b 100644 --- a/Help/prop_tgt/HIP_ARCHITECTURES.rst +++ b/Help/prop_tgt/HIP_ARCHITECTURES.rst @@ -5,8 +5,9 @@ HIP_ARCHITECTURES List of AMD GPU architectures to generate device code for. -An empty or false value (e.g. ``OFF``) defers architecture generation to compiler -defaults. +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. diff --git a/Help/variable/CMAKE_HIP_ARCHITECTURES.rst b/Help/variable/CMAKE_HIP_ARCHITECTURES.rst index 333f021..0cf0201 100644 --- a/Help/variable/CMAKE_HIP_ARCHITECTURES.rst +++ b/Help/variable/CMAKE_HIP_ARCHITECTURES.rst @@ -5,7 +5,7 @@ CMAKE_HIP_ARCHITECTURES Default value for :prop_tgt:`HIP_ARCHITECTURES` property of targets. -This is initialized to ``OFF``. +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/Modules/CMakeDetermineHIPCompiler.cmake b/Modules/CMakeDetermineHIPCompiler.cmake index 87885c0..ed0110a 100644 --- a/Modules/CMakeDetermineHIPCompiler.cmake +++ b/Modules/CMakeDetermineHIPCompiler.cmake @@ -85,7 +85,12 @@ if(MSVC_HIP_ARCHITECTURE_ID) endif() if(NOT DEFINED CMAKE_HIP_ARCHITECTURES) - set(CMAKE_HIP_ARCHITECTURES "OFF" CACHE STRING "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 diff --git a/Source/cmGeneratorTarget.cxx b/Source/cmGeneratorTarget.cxx index cb10117..f268c6c 100644 --- a/Source/cmGeneratorTarget.cxx +++ b/Source/cmGeneratorTarget.cxx @@ -3419,6 +3419,12 @@ 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; diff --git a/Tests/HIP/InferHipLang1/CMakeLists.txt b/Tests/HIP/InferHipLang1/CMakeLists.txt index f3c5421..63d77fd 100644 --- a/Tests/HIP/InferHipLang1/CMakeLists.txt +++ b/Tests/HIP/InferHipLang1/CMakeLists.txt @@ -10,4 +10,3 @@ target_compile_features(InterfaceWithHIP INTERFACE cxx_std_11) add_executable(HIPInferHipLang1 ) target_link_libraries(HIPInferHipLang1 PRIVATE InterfaceWithHIP) -set_property(TARGET HIPInferHipLang1 PROPERTY HIP_ARCHITECTURES gfx900) diff --git a/Tests/HIP/InferHipLang2/CMakeLists.txt b/Tests/HIP/InferHipLang2/CMakeLists.txt index a8a55b7..0e69de3 100644 --- a/Tests/HIP/InferHipLang2/CMakeLists.txt +++ b/Tests/HIP/InferHipLang2/CMakeLists.txt @@ -7,7 +7,6 @@ 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) -set_property(TARGET InterfaceWithHIP PROPERTY HIP_ARCHITECTURES gfx900) add_executable(HIPInferHipLang2 ) target_link_libraries(HIPInferHipLang2 PRIVATE InterfaceWithHIP) diff --git a/Tests/HIP/MathFunctions/CMakeLists.txt b/Tests/HIP/MathFunctions/CMakeLists.txt index 69d63dd..81e3ddb 100644 --- a/Tests/HIP/MathFunctions/CMakeLists.txt +++ b/Tests/HIP/MathFunctions/CMakeLists.txt @@ -13,7 +13,6 @@ project(MathFunctions HIP) # 3. This makes sure CMake properly links to all the built-in libraries # that hip needs that inject support for __half support # -set(CMAKE_HIP_ARCHITECTURES "gfx900") add_executable(HIPOnlyMathFunctions main.hip) target_compile_options(HIPOnlyMathFunctions PRIVATE -Werror) target_compile_features(HIPOnlyMathFunctions PRIVATE hip_std_14) diff --git a/Tests/HIP/MixedLanguage/CMakeLists.txt b/Tests/HIP/MixedLanguage/CMakeLists.txt index ce2ff89..4f6dd3b 100644 --- a/Tests/HIP/MixedLanguage/CMakeLists.txt +++ b/Tests/HIP/MixedLanguage/CMakeLists.txt @@ -3,7 +3,6 @@ project (MixedLanguage C CXX HIP) set(CMAKE_HIP_STANDARD 14) set(CMAKE_CXX_STANDARD 14) -set(CMAKE_HIP_ARCHITECTURES "gfx900") #Goal for this example: #make sure that we can build multiple languages into targets -- cgit v0.12