HIP: analyze output of hipcc to determine default GPU architecture

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.

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.

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

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;

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)

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)

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)

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