echo/CMake
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

HIP: Add support for NVIDIA GPUs

Browse Source 
Add support for using the CUDA Toolkit's NVCC to compile HIP code.

Fixes: #25143
This commit is contained in:
Brad King 2023-09-18 15:46:39 -04:00
parent 127b6fa06b
commit 18158bf81c
19 changed files with 221 additions and 27 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

5
Help/release/dev/hip-nvidia.rst
View File

@ -1,6 +1,9 @@
hip-nvidia
----------
* ``HIP`` language code may now be compiled for NVIDIA GPUs
using the NVIDIA CUDA Compiler (NVCC).
* The :variable:`CMAKE_HIP_PLATFORM` variable was added to specify
the GPU platform for which HIP language sources are to be compiled
(``amd``).
(``amd`` or ``nvidia``).

5
Help/variable/CMAKE_HIP_PLATFORM.rst
View File

@ -10,7 +10,10 @@ The value must be one of:
``amd``
AMD GPUs
If not specified, the default is ``amd``.
``nvidia``
NVIDIA GPUs
If not specified, a default is computed via ``hipconfig --platform``.
:variable:`CMAKE_HIP_ARCHITECTURES` entries are interpreted with
as architectures of the GPU platform.

2
Modules/CMakeCompilerIdDetection.cmake
View File

@ -103,7 +103,7 @@ function(compiler_id_detection outvar lang)
endif()
if("x${lang}" STREQUAL "xHIP")
set(ordered_compilers Clang)
set(ordered_compilers NVIDIA Clang)
endif()
if(CID_ID_DEFINE)

102
Modules/CMakeDetermineHIPCompiler.cmake
View File

@ -11,13 +11,21 @@ if( NOT ( ("${CMAKE_GENERATOR}" MATCHES "Make") OR
endif()
if(NOT CMAKE_HIP_PLATFORM)
set(CMAKE_HIP_PLATFORM "amd" CACHE STRING "HIP platform" FORCE)
execute_process(COMMAND hipconfig --platform
OUTPUT_VARIABLE _CMAKE_HIPCONFIG_PLATFORM OUTPUT_STRIP_TRAILING_WHITESPACE
RESULT_VARIABLE _CMAKE_HIPCONFIG_RESULT
)
if(_CMAKE_HIPCONFIG_RESULT EQUAL 0 AND _CMAKE_HIPCONFIG_PLATFORM MATCHES "^(nvidia|nvcc)$")
set(CMAKE_HIP_PLATFORM "nvidia" CACHE STRING "HIP platform" FORCE)
else()
set(CMAKE_HIP_PLATFORM "amd" CACHE STRING "HIP platform" FORCE)
endif()
endif()
if(NOT CMAKE_HIP_PLATFORM MATCHES "^(amd)$")
if(NOT CMAKE_HIP_PLATFORM MATCHES "^(amd|nvidia)$")
message(FATAL_ERROR
"The CMAKE_HIP_PLATFORM has unsupported value:\n"
" '${CMAKE_HIP_PLATFORM}'\n"
"It must be 'amd'."
"It must be 'amd' or 'nvidia'."
)
endif()
@ -44,7 +52,9 @@ if(NOT CMAKE_HIP_COMPILER)
# finally list compilers to try
if(NOT CMAKE_HIP_COMPILER_INIT)
if(CMAKE_HIP_PLATFORM STREQUAL "amd")
if(CMAKE_HIP_PLATFORM STREQUAL "nvidia")
set(CMAKE_HIP_COMPILER_LIST nvcc)
elseif(CMAKE_HIP_PLATFORM STREQUAL "amd")
set(CMAKE_HIP_COMPILER_LIST clang++)
# Look for the Clang coming with ROCm to support HIP.
@ -75,17 +85,55 @@ mark_as_advanced(CMAKE_HIP_COMPILER)
if(NOT CMAKE_HIP_COMPILER_ID_RUN)
set(CMAKE_HIP_COMPILER_ID_RUN 1)
# Try to identify the compiler.
include(${CMAKE_ROOT}/Modules/CMakeDetermineCompilerId.cmake)
# We determine the vendor to use the right flags for detection right away.
# The main compiler identification is still needed below to extract other information.
list(APPEND CMAKE_HIP_COMPILER_ID_VENDORS NVIDIA Clang)
set(CMAKE_HIP_COMPILER_ID_VENDOR_REGEX_NVIDIA "nvcc: NVIDIA \\(R\\) Cuda compiler driver")
set(CMAKE_HIP_COMPILER_ID_VENDOR_REGEX_Clang "(clang version)")
CMAKE_DETERMINE_COMPILER_ID_VENDOR(HIP "--version")
if(CMAKE_HIP_COMPILER_ID STREQUAL "NVIDIA")
# Find the CUDA toolkit to get:
# - CMAKE_HIP_COMPILER_CUDA_TOOLKIT_VERSION
# - CMAKE_HIP_COMPILER_CUDA_TOOLKIT_ROOT
# - CMAKE_HIP_COMPILER_CUDA_LIBRARY_ROOT
# We save them in CMakeHIPCompiler.cmake.
# Match arguments with cmake_cuda_architectures_all call.
include(Internal/CMakeCUDAFindToolkit)
cmake_cuda_find_toolkit(HIP CMAKE_HIP_COMPILER_CUDA_)
# If the user set CMAKE_HIP_ARCHITECTURES, validate its value.
include(Internal/CMakeCUDAArchitecturesValidate)
cmake_cuda_architectures_validate(HIP)
endif()
if(CMAKE_HIP_COMPILER_ID STREQUAL "Clang")
list(APPEND CMAKE_HIP_COMPILER_ID_TEST_FLAGS_FIRST "-v")
elseif(CMAKE_HIP_COMPILER_ID STREQUAL "NVIDIA")
# Tell nvcc to treat .hip files as CUDA sources.
list(APPEND CMAKE_HIP_COMPILER_ID_TEST_FLAGS_FIRST "-x cu -v")
endif()
# We perform compiler identification for a second time to extract implicit linking info.
# We need to unset the compiler ID otherwise CMAKE_DETERMINE_COMPILER_ID() doesn't work.
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)
if(CMAKE_HIP_COMPILER_ID STREQUAL "NVIDIA")
include(Internal/CMakeCUDAArchitecturesAll)
# From CMAKE_HIP_COMPILER_CUDA_TOOLKIT_VERSION and CMAKE_HIP_COMPILER_{ID,VERSION}, get:
# - CMAKE_HIP_ARCHITECTURES_ALL
# - CMAKE_HIP_ARCHITECTURES_ALL_MAJOR
# Match arguments with cmake_cuda_find_toolkit call.
cmake_cuda_architectures_all(HIP CMAKE_HIP_COMPILER_CUDA_)
endif()
_cmake_find_compiler_sysroot(HIP)
endif()
@ -116,7 +164,8 @@ if(NOT CMAKE_HIP_COMPILER_ROCM_ROOT)
message(FATAL_ERROR "Failed to find ROCm root directory.")
endif()
# Normally implicit link information is not detected until
# Normally implicit link information is not detected until ABI detection,
# but we need to populate CMAKE_HIP_LIBRARY_ARCHITECTURE to find hip-lang.
cmake_parse_implicit_link_info("${CMAKE_HIP_COMPILER_PRODUCED_OUTPUT}"
_CMAKE_HIP_COMPILER_ID_IMPLICIT_LIBS
_CMAKE_HIP_COMPILER_ID_IMPLICIT_DIRS
@ -177,6 +226,26 @@ include(CMakeFindBinUtils)
include(Compiler/${CMAKE_HIP_COMPILER_ID}-FindBinUtils OPTIONAL)
unset(_CMAKE_PROCESSING_LANGUAGE)
if(CMAKE_HIP_COMPILER_ID STREQUAL "Clang")
set(CMAKE_HIP_RUNTIME_LIBRARY_DEFAULT "SHARED")
elseif(CMAKE_HIP_COMPILER_ID STREQUAL "NVIDIA")
include(Internal/CMakeNVCCParseImplicitInfo)
# Parse CMAKE_HIP_COMPILER_PRODUCED_OUTPUT to get:
# - CMAKE_HIP_ARCHITECTURES_DEFAULT
# - CMAKE_HIP_HOST_IMPLICIT_LINK_DIRECTORIES
# - CMAKE_HIP_HOST_IMPLICIT_LINK_FRAMEWORK_DIRECTORIES
# - CMAKE_HIP_HOST_IMPLICIT_LINK_LIBRARIES
# - CMAKE_HIP_HOST_LINK_LAUNCHER
# - CMAKE_HIP_RUNTIME_LIBRARY_DEFAULT
# - CMAKE_HIP_CUDA_TOOLKIT_INCLUDE_DIRECTORIES
# Match arguments with cmake_nvcc_filter_implicit_info call in CMakeTestHIPCompiler.
cmake_nvcc_parse_implicit_info(HIP CMAKE_HIP_CUDA_)
include(Internal/CMakeCUDAFilterImplicitLibs)
# Filter out implicit link libraries that should not be passed unconditionally.
cmake_cuda_filter_implicit_libs(CMAKE_HIP_HOST_IMPLICIT_LINK_LIBRARIES)
endif()
if(CMAKE_HIP_COMPILER_SYSROOT)
string(CONCAT _SET_CMAKE_HIP_COMPILER_SYSROOT
"set(CMAKE_HIP_COMPILER_SYSROOT \"${CMAKE_HIP_COMPILER_SYSROOT}\")\n"
@ -197,7 +266,20 @@ if(MSVC_HIP_ARCHITECTURE_ID)
"set(MSVC_HIP_ARCHITECTURE_ID ${MSVC_HIP_ARCHITECTURE_ID})")
endif()
if(NOT DEFINED CMAKE_HIP_ARCHITECTURES)
if(CMAKE_HIP_COMPILER_ID STREQUAL "NVIDIA")
if(NOT "$ENV{CUDAARCHS}" STREQUAL "")
set(CMAKE_HIP_ARCHITECTURES "$ENV{CUDAARCHS}" CACHE STRING "CUDA architectures")
endif()
# If the user did not set CMAKE_HIP_ARCHITECTURES, use the compiler's default.
if("${CMAKE_HIP_ARCHITECTURES}" STREQUAL "")
set(CMAKE_HIP_ARCHITECTURES "${CMAKE_HIP_ARCHITECTURES_DEFAULT}" CACHE STRING "HIP architectures" FORCE)
if(NOT CMAKE_HIP_ARCHITECTURES)
message(FATAL_ERROR "Failed to detect a default HIP architecture.\n\nCompiler output:\n${CMAKE_HIP_COMPILER_PRODUCED_OUTPUT}")
endif()
endif()
unset(CMAKE_HIP_ARCHITECTURES_DEFAULT)
elseif(NOT DEFINED CMAKE_HIP_ARCHITECTURES)
# Use 'rocm_agent_enumerator' to get the current GPU architecture.
set(_CMAKE_HIP_ARCHITECTURES)
find_program(_CMAKE_HIP_ROCM_AGENT_ENUMERATOR

19
Modules/CMakeHIPCompiler.cmake.in
View File

@ -1,4 +1,6 @@
set(CMAKE_HIP_COMPILER "@CMAKE_HIP_COMPILER@")
set(CMAKE_HIP_HOST_COMPILER "@CMAKE_HIP_HOST_COMPILER@")
set(CMAKE_HIP_HOST_LINK_LAUNCHER "@CMAKE_HIP_HOST_LINK_LAUNCHER@")
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@")
@ -45,14 +47,27 @@ 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_COMPILER_CUDA_TOOLKIT_ROOT "@CMAKE_HIP_COMPILER_CUDA_TOOLKIT_ROOT@")
set(CMAKE_HIP_COMPILER_CUDA_TOOLKIT_LIBRARY_ROOT "@CMAKE_HIP_COMPILER_CUDA_TOOLKIT_LIBRARY_ROOT@")
set(CMAKE_HIP_COMPILER_CUDA_TOOLKIT_VERSION "@CMAKE_HIP_COMPILER_CUDA_TOOLKIT_VERSION@")
set(CMAKE_HIP_COMPILER_CUDA_LIBRARY_ROOT "@CMAKE_HIP_COMPILER_CUDA_LIBRARY_ROOT@")
set(CMAKE_HIP_ARCHITECTURES_ALL "@CMAKE_HIP_ARCHITECTURES_ALL@")
set(CMAKE_HIP_ARCHITECTURES_ALL_MAJOR "@CMAKE_HIP_ARCHITECTURES_ALL_MAJOR@")
set(CMAKE_HIP_ARCHITECTURES_NATIVE "@CMAKE_HIP_ARCHITECTURES_NATIVE@")
set(CMAKE_HIP_CUDA_TOOLKIT_INCLUDE_DIRECTORIES "@CMAKE_HIP_CUDA_TOOLKIT_INCLUDE_DIRECTORIES@")
set(CMAKE_HIP_HOST_IMPLICIT_LINK_LIBRARIES "@CMAKE_HIP_HOST_IMPLICIT_LINK_LIBRARIES@")
set(CMAKE_HIP_HOST_IMPLICIT_LINK_DIRECTORIES "@CMAKE_HIP_HOST_IMPLICIT_LINK_DIRECTORIES@")
set(CMAKE_HIP_HOST_IMPLICIT_LINK_FRAMEWORK_DIRECTORIES "@CMAKE_HIP_HOST_IMPLICIT_LINK_FRAMEWORK_DIRECTORIES@")
set(CMAKE_HIP_IMPLICIT_INCLUDE_DIRECTORIES "@CMAKE_HIP_IMPLICIT_INCLUDE_DIRECTORIES@")
set(CMAKE_HIP_IMPLICIT_LINK_LIBRARIES "@CMAKE_HIP_IMPLICIT_LINK_LIBRARIES@")
set(CMAKE_HIP_IMPLICIT_LINK_DIRECTORIES "@CMAKE_HIP_IMPLICIT_LINK_DIRECTORIES@")
set(CMAKE_HIP_IMPLICIT_LINK_FRAMEWORK_DIRECTORIES "@CMAKE_HIP_IMPLICIT_LINK_FRAMEWORK_DIRECTORIES@")
set(CMAKE_HIP_RUNTIME_LIBRARY_DEFAULT "SHARED")
set(CMAKE_HIP_RUNTIME_LIBRARY_DEFAULT "@CMAKE_HIP_RUNTIME_LIBRARY_DEFAULT@")
set(CMAKE_AR "@CMAKE_AR@")
set(CMAKE_HIP_COMPILER_AR "@CMAKE_HIP_COMPILER_AR@")

18
Modules/CMakeHIPCompilerABI.hip
View File

@ -1,9 +1,13 @@
#ifndef __HIP__
#if !defined(__HIP__) && !defined(__NVCC__)
# error "A C or C++ compiler has been selected for HIP"
#endif
#include "CMakeCompilerABI.h"
#if defined(__NVCC__)
# include "CMakeCompilerCUDAArch.h"
#endif
int main(int argc, char* argv[])
{
int require = 0;
@ -11,6 +15,16 @@ int main(int argc, char* argv[])
#if defined(ABI_ID)
require += info_abi[argc];
#endif
(void)argv;
static_cast<void>(argv);
#if defined(__NVCC__)
if (!cmakeCompilerCUDAArch()) {
// Convince the compiler that the non-zero return value depends
// on the info strings so they are not optimized out.
return require ? -1 : 1;
}
return 0;
#else
return require;
#endif
}

2
Modules/CMakeHIPCompilerId.hip.in
View File

@ -1,4 +1,4 @@
#ifndef __HIP__
#if !defined(__HIP__) && !defined(__NVCC__)
# error "A C or C++ compiler has been selected for HIP"
#endif

15
Modules/CMakeHIPInformation.cmake
View File

@ -8,6 +8,19 @@ else()
endif()
set(CMAKE_INCLUDE_FLAG_HIP "-I")
# Set implicit links early so compiler-specific modules can use them.
set(__IMPLICIT_LINKS)
foreach(dir ${CMAKE_HIP_HOST_IMPLICIT_LINK_DIRECTORIES})
string(APPEND __IMPLICIT_LINKS " -L\"${dir}\"")
endforeach()
foreach(lib ${CMAKE_HIP_HOST_IMPLICIT_LINK_LIBRARIES})
if(${lib} MATCHES "/")
string(APPEND __IMPLICIT_LINKS " \"${lib}\"")
else()
string(APPEND __IMPLICIT_LINKS " -l${lib}")
endif()
endforeach()
# Load compiler-specific information.
if(CMAKE_HIP_COMPILER_ID)
include(Compiler/${CMAKE_HIP_COMPILER_ID}-HIP OPTIONAL)
@ -129,7 +142,7 @@ endif()
# compile a HIP file into an object file
if(NOT CMAKE_HIP_COMPILE_OBJECT)
set(CMAKE_HIP_COMPILE_OBJECT
"<CMAKE_HIP_COMPILER> <DEFINES> <INCLUDES> <FLAGS> -o <OBJECT> ${_CMAKE_COMPILE_AS_HIP_FLAG} -c <SOURCE>")
"<CMAKE_HIP_COMPILER> ${_CMAKE_HIP_EXTRA_FLAGS} <DEFINES> <INCLUDES> <FLAGS> -o <OBJECT> ${_CMAKE_COMPILE_AS_HIP_FLAG} -c <SOURCE>")
endif()
# compile a cu file into an executable

24
Modules/CMakeTestHIPCompiler.cmake
View File

@ -10,7 +10,10 @@ if(CMAKE_HIP_COMPILER_FORCED)
endif()
set(__CMAKE_HIP_FLAGS "${CMAKE_HIP_FLAGS}")
string(APPEND CMAKE_HIP_FLAGS " --cuda-host-only")
if(CMAKE_HIP_COMPILER_ID STREQUAL "Clang")
string(APPEND CMAKE_HIP_FLAGS " --cuda-host-only")
endif()
include(CMakeTestCompilerCommon)
@ -31,6 +34,13 @@ 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")
if(CMAKE_HIP_COMPILER_ID STREQUAL "NVIDIA")
include(Internal/CMakeCUDAArchitecturesNative)
# Run the test binary to get:
# - CMAKE_HIP_ARCHITECTURES_NATIVE
cmake_cuda_architectures_native(HIP)
endif()
endif()
# This file is used by EnableLanguage in cmGlobalGenerator to
@ -42,7 +52,7 @@ if(NOT CMAKE_HIP_COMPILER_WORKS)
PrintTestCompilerStatus("HIP")
__TestCompiler_setTryCompileTargetType()
string(CONCAT __TestCompiler_testHIPCompilerSource
"#ifndef __HIP__\n"
"#if !defined(__HIP__) && !defined(__NVCC__)\n"
"# error \"The CMAKE_HIP_COMPILER is set to a C/CXX compiler\"\n"
"#endif\n"
"int main(){return 0;}\n")
@ -76,6 +86,16 @@ unset(__CMAKE_HIP_FLAGS)
include(${CMAKE_ROOT}/Modules/CMakeDetermineCompileFeatures.cmake)
CMAKE_DETERMINE_COMPILE_FEATURES(HIP)
if(CMAKE_HIP_COMPILER_ID STREQUAL "NVIDIA")
include(Internal/CMakeNVCCFilterImplicitInfo)
# Match arguments with cmake_nvcc_parse_implicit_info call in CMakeDetermineHIPCompiler.
cmake_nvcc_filter_implicit_info(HIP CMAKE_HIP_CUDA_)
include(Internal/CMakeCUDAFilterImplicitLibs)
# Filter out implicit link libraries that should not be passed unconditionally.
cmake_cuda_filter_implicit_libs(CMAKE_HIP_IMPLICIT_LINK_LIBRARIES)
endif()
# Re-configure to save learned information.
configure_file(
${CMAKE_ROOT}/Modules/CMakeHIPCompiler.cmake.in

16
Modules/Compiler/NVIDIA-HIP.cmake Normal file
View File

@ -0,0 +1,16 @@
include(Compiler/NVIDIA)
__compiler_nvidia_cxx_standards(HIP)
__compiler_nvidia_cuda_flags(HIP)
# The ROCm hip-lang package does not work for nvcc,
# so provide a minimal one ourselves.
add_library(hip-lang::device INTERFACE IMPORTED)
set(_CMAKE_HIP_DEVICE_RUNTIME_TARGET hip-lang::device)
set(CMAKE_HIP_STANDARD_INCLUDE_DIRECTORIES "${CMAKE_HIP_COMPILER_ROCM_ROOT}/include")
set(CMAKE_HIP_LINK_EXECUTABLE
"<CMAKE_HIP_HOST_LINK_LAUNCHER> <LINK_FLAGS> <OBJECTS> -o <TARGET> <LINK_LIBRARIES>${__IMPLICIT_LINKS}")
set(CMAKE_HIP_CREATE_SHARED_LIBRARY
"<CMAKE_HIP_HOST_LINK_LAUNCHER> <CMAKE_SHARED_LIBRARY_HIP_FLAGS> <LINK_FLAGS> <CMAKE_SHARED_LIBRARY_CREATE_HIP_FLAGS> <SONAME_FLAG><TARGET_SONAME> -o <TARGET> <OBJECTS> <LINK_LIBRARIES>${__IMPLICIT_LINKS}")
set(CMAKE_HIP_CREATE_SHARED_MODULE "${CMAKE_HIP_CREATE_SHARED_LIBRARY}")

11
Source/cmGeneratorTarget.cxx
View File

@ -3579,9 +3579,11 @@ void cmGeneratorTarget::AddISPCTargetFlags(std::string& flags) const
}
}
void cmGeneratorTarget::AddHIPArchitectureFlags(std::string& flags) const
void cmGeneratorTarget::AddHIPArchitectureFlags(cmBuildStep compileOrLink,
std::string const& config,
std::string& flags) const
{
const std::string& arch = this->GetSafeProperty("HIP_ARCHITECTURES");
std::string arch = this->GetSafeProperty("HIP_ARCHITECTURES");
if (arch.empty()) {
this->Makefile->IssueMessage(MessageType::FATAL_ERROR,
@ -3594,6 +3596,11 @@ void cmGeneratorTarget::AddHIPArchitectureFlags(std::string& flags) const
return;
}
if (this->Makefile->GetSafeDefinition("CMAKE_HIP_PLATFORM") == "nvidia") {
return this->AddCUDAArchitectureFlagsImpl(compileOrLink, config, "HIP",
std::move(arch), flags);
}
cmList options(arch);
for (std::string& option : options) {

4
Source/cmGeneratorTarget.h
View File

@ -504,7 +504,9 @@ public:
std::string& flags) const;
void AddCUDAToolkitFlags(std::string& flags) const;
void AddHIPArchitectureFlags(std::string& flags) const;
void AddHIPArchitectureFlags(cmBuildStep compileOrLink,
std::string const& config,
std::string& flags) const;
void AddISPCTargetFlags(std::string& flags) const;

4
Source/cmLocalGenerator.cxx
View File

@ -87,6 +87,8 @@ static auto ruleReplaceVars = { "CMAKE_${LANG}_COMPILER",
"CMAKE_TAPI",
"CMAKE_CUDA_HOST_COMPILER",
"CMAKE_CUDA_HOST_LINK_LAUNCHER",
"CMAKE_HIP_HOST_COMPILER",
"CMAKE_HIP_HOST_LINK_LAUNCHER",
"CMAKE_CL_SHOWINCLUDES_PREFIX" };
cmLocalGenerator::cmLocalGenerator(cmGlobalGenerator* gg, cmMakefile* makefile)
@ -2058,7 +2060,7 @@ void cmLocalGenerator::AddLanguageFlags(std::string& flags,
this->Makefile->GetSafeDefinition("CMAKE_CXX_SIMULATE_ID");
}
} else if (lang == "HIP") {
target->AddHIPArchitectureFlags(flags);
target->AddHIPArchitectureFlags(compileOrLink, config, flags);
}
// Add VFS Overlay for Clang compilers

2
Tests/HIP/ArchitectureOff/CMakeLists.txt
View File

@ -7,6 +7,8 @@ set(CMAKE_HIP_ARCHITECTURES OFF)
# Pass our own architecture flags instead.
if(CMAKE_HIP_PLATFORM STREQUAL "amd")
string(APPEND CMAKE_HIP_FLAGS " --offload-arch=gfx908")
elseif(CMAKE_HIP_PLATFORM STREQUAL "nvidia")
string(APPEND CMAKE_HIP_FLAGS " -arch=sm_52")
endif()
add_executable(HIPOnlyArchitectureOff main.hip)

5
Tests/HIP/CMakeLists.txt
View File

@ -9,7 +9,10 @@ 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)
if(CMake_TEST_HIP STREQUAL "amd")
# The NVIDIA CUDA compiler cannot handle device lambda markup.
add_hip_test_macro(HIP.MathFunctions HIPOnlyMathFunctions)
endif()
add_hip_test_macro(HIP.MixedLanguage HIPMixedLanguage)
add_hip_test_macro(HIP.TryCompile HIPOnlyTryCompile)
add_hip_test_macro(HIP.WithDefs HIPOnlyWithDefs)

2
Tests/HIP/CompileFlags/CMakeLists.txt
View File

@ -5,6 +5,8 @@ add_executable(HIPOnlyCompileFlags main.hip)
if(CMAKE_HIP_PLATFORM STREQUAL "amd")
set(hip_archs gfx803)
elseif(CMAKE_HIP_PLATFORM STREQUAL "nvidia")
set(hip_archs 52)
endif()
set_property(TARGET HIPOnlyCompileFlags PROPERTY HIP_ARCHITECTURES ${hip_archs})

6
Tests/HIP/MathFunctions/CMakeLists.txt
View File

@ -14,5 +14,9 @@ project(MathFunctions HIP)
# that hip needs that inject support for __half support
#
add_executable(HIPOnlyMathFunctions main.hip)
target_compile_options(HIPOnlyMathFunctions PRIVATE -Werror)
if(CMAKE_HIP_COMPILER_ID STREQUAL "NVIDIA")
target_compile_options(HIPOnlyMathFunctions PRIVATE "SHELL:-Werror all-warnings")
elseif(CMAKE_HIP_COMPILER_ID STREQUAL "Clang")
target_compile_options(HIPOnlyMathFunctions PRIVATE "-Werror")
endif()
target_compile_features(HIPOnlyMathFunctions PRIVATE hip_std_14)

2
Tests/HIP/TryCompile/CMakeLists.txt
View File

@ -7,6 +7,8 @@ set(CMAKE_HIP_STANDARD 14)
if(CMAKE_HIP_PLATFORM STREQUAL "amd")
set(CMAKE_HIP_ARCHITECTURES gfx803 gfx900)
elseif(CMAKE_HIP_PLATFORM STREQUAL "nvidia")
set(CMAKE_HIP_ARCHITECTURES 52)
endif()
set(CMAKE_TRY_COMPILE_TARGET_TYPE STATIC_LIBRARY)

4
Tests/HIP/WithDefs/main.hip.cpp
View File

@ -51,6 +51,10 @@ static __global__ void DetermineIfValidHIPDevice()
# undef PACKED_DEFINE
# define PACKED_DEFINE
#endif
#ifdef __NVCC__
# undef PACKED_DEFINE
# define PACKED_DEFINE
#endif
struct PACKED_DEFINE result_type
{
bool valid;