From 1380a9a672a0302a39de0f9faf78fffafb318119 Mon Sep 17 00:00:00 2001 From: Jens Glaser Date: Mon, 13 Jun 2022 17:27:52 -0400 Subject: [PATCH 01/15] Initial changes to compile on AMD --- CMake/HIP/FindHIP.cmake | 577 ------------------------- CMake/HIP/FindHIP/run_hipcc.cmake | 168 ------- CMake/HIP/FindHIP/run_make2cmake.cmake | 50 --- CMake/hoomd/FindCUDALibs.cmake | 126 +----- CMake/hoomd/HOOMDHIPSetup.cmake | 164 +++---- CMake/hoomd/hipcc.cmake | 51 --- CMakeLists.txt | 7 +- hoomd/CMakeLists.txt | 6 +- hoomd/SnapshotSystemData.cc | 2 +- hoomd/md/BondTablePotential.cc | 4 +- hoomd/md/EvaluatorTersoff.h | 8 +- hoomd/md/ForceComposite.h | 14 +- hoomd/md/OPLSDihedralForceCompute.cc | 32 +- hoomd/md/TwoStepRATTLEBDGPU.h | 4 + 14 files changed, 106 insertions(+), 1107 deletions(-) delete mode 100644 CMake/HIP/FindHIP.cmake delete mode 100644 CMake/HIP/FindHIP/run_hipcc.cmake delete mode 100644 CMake/HIP/FindHIP/run_make2cmake.cmake delete mode 100644 CMake/hoomd/hipcc.cmake diff --git a/CMake/HIP/FindHIP.cmake b/CMake/HIP/FindHIP.cmake deleted file mode 100644 index 16bc8f6ee1..0000000000 --- a/CMake/HIP/FindHIP.cmake +++ /dev/null @@ -1,577 +0,0 @@ -############################################################################### -# FindHIP.cmake -############################################################################### - -############################################################################### -# SET: Variable defaults -############################################################################### -# User defined flags -set(HIP_HIPCC_FLAGS "" CACHE STRING "Semicolon delimited flags for HIPCC") -set(HIP_HCC_FLAGS "" CACHE STRING "Semicolon delimited flags for HCC") -set(HIP_NVCC_FLAGS "" CACHE STRING "Semicolon delimted flags for NVCC") -mark_as_advanced(HIP_HIPCC_FLAGS HIP_HCC_FLAGS HIP_NVCC_FLAGS) -set(_hip_configuration_types ${CMAKE_CONFIGURATION_TYPES} ${CMAKE_BUILD_TYPE} Debug MinSizeRel Release RelWithDebInfo) -list(REMOVE_DUPLICATES _hip_configuration_types) -foreach(config ${_hip_configuration_types}) - string(TOUPPER ${config} config_upper) - set(HIP_HIPCC_FLAGS_${config_upper} "" CACHE STRING "Semicolon delimited flags for HIPCC") - set(HIP_HCC_FLAGS_${config_upper} "" CACHE STRING "Semicolon delimited flags for HCC") - set(HIP_NVCC_FLAGS_${config_upper} "" CACHE STRING "Semicolon delimited flags for NVCC") - mark_as_advanced(HIP_HIPCC_FLAGS_${config_upper} HIP_HCC_FLAGS_${config_upper} HIP_NVCC_FLAGS_${config_upper}) -endforeach() -option(HIP_HOST_COMPILATION_CPP "Host code compilation mode" ON) -option(HIP_VERBOSE_BUILD "Print out the commands run while compiling the HIP source file. With the Makefile generator this defaults to VERBOSE variable specified on the command line, but can be forced on with this option." OFF) -mark_as_advanced(HIP_HOST_COMPILATION_CPP) - -############################################################################### -# Set HIP CMAKE Flags -############################################################################### -# Copy the invocation styles from CXX to HIP -set(CMAKE_HIP_ARCHIVE_CREATE ${CMAKE_CXX_ARCHIVE_CREATE}) -set(CMAKE_HIP_ARCHIVE_APPEND ${CMAKE_CXX_ARCHIVE_APPEND}) -set(CMAKE_HIP_ARCHIVE_FINISH ${CMAKE_CXX_ARCHIVE_FINISH}) -set(CMAKE_SHARED_LIBRARY_SONAME_HIP_FLAG ${CMAKE_SHARED_LIBRARY_SONAME_CXX_FLAG}) -set(CMAKE_SHARED_LIBRARY_CREATE_HIP_FLAGS ${CMAKE_SHARED_LIBRARY_CREATE_CXX_FLAGS}) -set(CMAKE_SHARED_LIBRARY_HIP_FLAGS ${CMAKE_SHARED_LIBRARY_CXX_FLAGS}) -#set(CMAKE_SHARED_LIBRARY_LINK_HIP_FLAGS ${CMAKE_SHARED_LIBRARY_LINK_CXX_FLAGS}) -set(CMAKE_SHARED_LIBRARY_RUNTIME_HIP_FLAG ${CMAKE_SHARED_LIBRARY_RUNTIME_CXX_FLAG}) -set(CMAKE_SHARED_LIBRARY_RUNTIME_HIP_FLAG_SEP ${CMAKE_SHARED_LIBRARY_RUNTIME_CXX_FLAG_SEP}) -set(CMAKE_SHARED_LIBRARY_LINK_STATIC_HIP_FLAGS ${CMAKE_SHARED_LIBRARY_LINK_STATIC_CXX_FLAGS}) -set(CMAKE_SHARED_LIBRARY_LINK_DYNAMIC_HIP_FLAGS ${CMAKE_SHARED_LIBRARY_LINK_DYNAMIC_CXX_FLAGS}) - -# Set the CMake Flags to use the HCC Compilier. -set(CMAKE_HIP_CREATE_SHARED_LIBRARY "${HIP_HIPCC_CMAKE_LINKER_HELPER} ${HCC_PATH} -o ") -set(CMAKE_HIP_CREATE_SHARED_MODULE "${HIP_HIPCC_CMAKE_LINKER_HELPER} ${HCC_PATH} -o -shared" ) -set(CMAKE_HIP_LINK_EXECUTABLE "${HIP_HIPCC_CMAKE_LINKER_HELPER} ${HCC_PATH} -o ") - -############################################################################### -# FIND: HIP and associated helper binaries -############################################################################### -# HIP is supported on Linux only -if(UNIX AND NOT APPLE AND NOT CYGWIN) - # Search for HIP installation - if(NOT HIP_ROOT_DIR) - # Search in user specified path first - find_path( - HIP_ROOT_DIR - NAMES hipconfig - PATHS - ENV ROCM_PATH - ENV HIP_PATH - PATH_SUFFIXES bin - DOC "HIP installed location" - NO_DEFAULT_PATH - ) - # Now search in default path - find_path( - HIP_ROOT_DIR - NAMES hipconfig - PATHS - /opt/rocm - /opt/rocm/hip - PATH_SUFFIXES bin - DOC "HIP installed location" - ) - - # Check if we found HIP installation - if(HIP_ROOT_DIR) - # If so, fix the path - string(REGEX REPLACE "[/\\\\]?bin[64]*[/\\\\]?$" "" HIP_ROOT_DIR ${HIP_ROOT_DIR}) - # And push it back to the cache - set(HIP_ROOT_DIR ${HIP_ROOT_DIR} CACHE PATH "HIP installed location" FORCE) - endif() - if(NOT EXISTS ${HIP_ROOT_DIR}) - if(HIP_FIND_REQUIRED) - message(FATAL_ERROR "Specify HIP_ROOT_DIR") - endif() - endif() - endif() - - # Find HIPCC executable - find_program( - HIP_HIPCC_EXECUTABLE - NAMES hipcc - PATHS - "${HIP_ROOT_DIR}" - ENV ROCM_PATH - ENV HIP_PATH - /opt/rocm - /opt/rocm/hip - PATH_SUFFIXES bin - NO_DEFAULT_PATH - ) - if(NOT HIP_HIPCC_EXECUTABLE) - # Now search in default paths - find_program(HIP_HIPCC_EXECUTABLE hipcc) - endif() - mark_as_advanced(HIP_HIPCC_EXECUTABLE) - - # Find HIPCONFIG executable - find_program( - HIP_HIPCONFIG_EXECUTABLE - NAMES hipconfig - PATHS - "${HIP_ROOT_DIR}" - ENV ROCM_PATH - ENV HIP_PATH - /opt/rocm - /opt/rocm/hip - PATH_SUFFIXES bin - NO_DEFAULT_PATH - ) - if(NOT HIP_HIPCONFIG_EXECUTABLE) - # Now search in default paths - find_program(HIP_HIPCONFIG_EXECUTABLE hipconfig) - endif() - mark_as_advanced(HIP_HIPCONFIG_EXECUTABLE) - - # Find HIPCC_CMAKE_LINKER_HELPER executable - find_program( - HIP_HIPCC_CMAKE_LINKER_HELPER - NAMES hipcc_cmake_linker_helper - PATHS - "${HIP_ROOT_DIR}" - ENV ROCM_PATH - ENV HIP_PATH - /opt/rocm - /opt/rocm/hip - PATH_SUFFIXES bin - NO_DEFAULT_PATH - ) - if(NOT HIP_HIPCC_CMAKE_LINKER_HELPER) - # Now search in default paths - find_program(HIP_HIPCC_CMAKE_LINKER_HELPER hipcc_cmake_linker_helper) - endif() - mark_as_advanced(HIP_HIPCC_CMAKE_LINKER_HELPER) - - if(HIP_HIPCONFIG_EXECUTABLE AND NOT HIP_VERSION) - # Compute the version - execute_process( - COMMAND ${HIP_HIPCONFIG_EXECUTABLE} --version - OUTPUT_VARIABLE _hip_version - ERROR_VARIABLE _hip_error - OUTPUT_STRIP_TRAILING_WHITESPACE - ERROR_STRIP_TRAILING_WHITESPACE - ) - if(NOT _hip_error) - set(HIP_VERSION ${_hip_version} CACHE STRING "Version of HIP as computed from hipcc") - else() - set(HIP_VERSION "0.0.0" CACHE STRING "Version of HIP as computed by FindHIP()") - endif() - mark_as_advanced(HIP_VERSION) - endif() - if(HIP_VERSION) - string(REPLACE "." ";" _hip_version_list "${HIP_VERSION}") - list(GET _hip_version_list 0 HIP_VERSION_MAJOR) - list(GET _hip_version_list 1 HIP_VERSION_MINOR) - list(GET _hip_version_list 2 HIP_VERSION_PATCH) - set(HIP_VERSION_STRING "${HIP_VERSION}") - endif() - - if(HIP_HIPCONFIG_EXECUTABLE AND NOT HIP_PLATFORM) - # Compute the platform - execute_process( - COMMAND ${HIP_HIPCONFIG_EXECUTABLE} --platform - OUTPUT_VARIABLE _hip_platform - OUTPUT_STRIP_TRAILING_WHITESPACE - ) - set(HIP_PLATFORM ${_hip_platform} CACHE STRING "HIP platform as computed by hipconfig") - mark_as_advanced(HIP_PLATFORM) - endif() -endif() - -include(FindPackageHandleStandardArgs) -find_package_handle_standard_args( - HIP - REQUIRED_VARS - HIP_ROOT_DIR - HIP_HIPCC_EXECUTABLE - HIP_HIPCONFIG_EXECUTABLE - HIP_PLATFORM - VERSION_VAR HIP_VERSION - ) - -############################################################################### -# MACRO: Locate helper files -############################################################################### -macro(HIP_FIND_HELPER_FILE _name _extension) - set(_hip_full_name "${_name}.${_extension}") - get_filename_component(CMAKE_CURRENT_LIST_DIR "${CMAKE_CURRENT_LIST_FILE}" PATH) - set(HIP_${_name} "${CMAKE_CURRENT_LIST_DIR}/FindHIP/${_hip_full_name}") - if(NOT EXISTS "${HIP_${_name}}") - set(error_message "${_hip_full_name} not found in ${CMAKE_CURRENT_LIST_DIR}/FindHIP") - if(HIP_FIND_REQUIRED) - message(FATAL_ERROR "${error_message}") - else() - if(NOT HIP_FIND_QUIETLY) - message(STATUS "${error_message}") - endif() - endif() - endif() - # Set this variable as internal, so the user isn't bugged with it. - set(HIP_${_name} ${HIP_${_name}} CACHE INTERNAL "Location of ${_full_name}" FORCE) -endmacro() - -############################################################################### -hip_find_helper_file(run_make2cmake cmake) -hip_find_helper_file(run_hipcc cmake) -############################################################################### - -############################################################################### -# MACRO: Reset compiler flags -############################################################################### -macro(HIP_RESET_FLAGS) - unset(HIP_HIPCC_FLAGS) - unset(HIP_HCC_FLAGS) - unset(HIP_NVCC_FLAGS) - foreach(config ${_hip_configuration_types}) - string(TOUPPER ${config} config_upper) - unset(HIP_HIPCC_FLAGS_${config_upper}) - unset(HIP_HCC_FLAGS_${config_upper}) - unset(HIP_NVCC_FLAGS_${config_upper}) - endforeach() -endmacro() - -############################################################################### -# MACRO: Separate the options from the sources -############################################################################### -macro(HIP_GET_SOURCES_AND_OPTIONS _sources _cmake_options _hipcc_options _hcc_options _nvcc_options) - set(${_sources}) - set(${_cmake_options}) - set(${_hipcc_options}) - set(${_hcc_options}) - set(${_nvcc_options}) - set(_hipcc_found_options FALSE) - set(_hcc_found_options FALSE) - set(_nvcc_found_options FALSE) - foreach(arg ${ARGN}) - if("x${arg}" STREQUAL "xHIPCC_OPTIONS") - set(_hipcc_found_options TRUE) - set(_hcc_found_options FALSE) - set(_nvcc_found_options FALSE) - elseif("x${arg}" STREQUAL "xHCC_OPTIONS") - set(_hipcc_found_options FALSE) - set(_hcc_found_options TRUE) - set(_nvcc_found_options FALSE) - elseif("x${arg}" STREQUAL "xNVCC_OPTIONS") - set(_hipcc_found_options FALSE) - set(_hcc_found_options FALSE) - set(_nvcc_found_options TRUE) - elseif( - "x${arg}" STREQUAL "xEXCLUDE_FROM_ALL" OR - "x${arg}" STREQUAL "xSTATIC" OR - "x${arg}" STREQUAL "xSHARED" OR - "x${arg}" STREQUAL "xMODULE" - ) - list(APPEND ${_cmake_options} ${arg}) - else() - if(_hipcc_found_options) - list(APPEND ${_hipcc_options} ${arg}) - elseif(_hcc_found_options) - list(APPEND ${_hcc_options} ${arg}) - elseif(_nvcc_found_options) - list(APPEND ${_nvcc_options} ${arg}) - else() - # Assume this is a file - list(APPEND ${_sources} ${arg}) - endif() - endif() - endforeach() -endmacro() - -############################################################################### -# MACRO: Add include directories to pass to the hipcc command -############################################################################### -set(HIP_HIPCC_INCLUDE_ARGS_USER "") -macro(HIP_INCLUDE_DIRECTORIES) - foreach(dir ${ARGN}) - list(APPEND HIP_HIPCC_INCLUDE_ARGS_USER $<$:-I${dir}>) - endforeach() -endmacro() - -############################################################################### -# FUNCTION: Helper to avoid clashes of files with the same basename but different paths -############################################################################### -function(HIP_COMPUTE_BUILD_PATH path build_path) - # Convert to cmake style paths - file(TO_CMAKE_PATH "${path}" bpath) - if(IS_ABSOLUTE "${bpath}") - string(FIND "${bpath}" "${CMAKE_CURRENT_BINARY_DIR}" _binary_dir_pos) - if(_binary_dir_pos EQUAL 0) - file(RELATIVE_PATH bpath "${CMAKE_CURRENT_BINARY_DIR}" "${bpath}") - else() - file(RELATIVE_PATH bpath "${CMAKE_CURRENT_SOURCE_DIR}" "${bpath}") - endif() - endif() - - # Remove leading / - string(REGEX REPLACE "^[/]+" "" bpath "${bpath}") - # Avoid absolute paths by removing ':' - string(REPLACE ":" "_" bpath "${bpath}") - # Avoid relative paths that go up the tree - string(REPLACE "../" "__/" bpath "${bpath}") - # Avoid spaces - string(REPLACE " " "_" bpath "${bpath}") - # Strip off the filename - get_filename_component(bpath "${bpath}" PATH) - - set(${build_path} "${bpath}" PARENT_SCOPE) -endfunction() - -############################################################################### -# MACRO: Parse OPTIONS from ARGN & set variables prefixed by _option_prefix -############################################################################### -macro(HIP_PARSE_HIPCC_OPTIONS _option_prefix) - set(_hip_found_config) - foreach(arg ${ARGN}) - # Determine if we are dealing with a per-configuration flag - foreach(config ${_hip_configuration_types}) - string(TOUPPER ${config} config_upper) - if(arg STREQUAL "${config_upper}") - set(_hip_found_config _${arg}) - # Clear arg to prevent it from being processed anymore - set(arg) - endif() - endforeach() - if(arg) - list(APPEND ${_option_prefix}${_hip_found_config} "${arg}") - endif() - endforeach() -endmacro() - -############################################################################### -# MACRO: Try and include dependency file if it exists -############################################################################### -macro(HIP_INCLUDE_HIPCC_DEPENDENCIES dependency_file) - set(HIP_HIPCC_DEPEND) - set(HIP_HIPCC_DEPEND_REGENERATE FALSE) - - # Create the dependency file if it doesn't exist - if(NOT EXISTS ${dependency_file}) - file(WRITE ${dependency_file} "# Generated by: FindHIP.cmake. Do not edit.\n") - endif() - # Include the dependency file - include(${dependency_file}) - - # Verify the existence of all the included files - if(HIP_HIPCC_DEPEND) - foreach(f ${HIP_HIPCC_DEPEND}) - if(NOT EXISTS ${f}) - # If they aren't there, regenerate the file again - set(HIP_HIPCC_DEPEND_REGENERATE TRUE) - endif() - endforeach() - else() - # No dependencies, so regenerate the file - set(HIP_HIPCC_DEPEND_REGENERATE TRUE) - endif() - - # Regenerate the dependency file if needed - if(HIP_HIPCC_DEPEND_REGENERATE) - set(HIP_HIPCC_DEPEND ${dependency_file}) - file(WRITE ${dependency_file} "# Generated by: FindHIP.cmake. Do not edit.\n") - endif() -endmacro() - -############################################################################### -# MACRO: Prepare cmake commands for the target -############################################################################### -macro(HIP_PREPARE_TARGET_COMMANDS _target _format _generated_files _source_files) - set(_hip_flags "") - string(TOUPPER "${CMAKE_BUILD_TYPE}" _hip_build_configuration) - if(HIP_HOST_COMPILATION_CPP) - set(HIP_C_OR_CXX CXX) - else() - set(HIP_C_OR_CXX C) - endif() - set(generated_extension ${CMAKE_${HIP_C_OR_CXX}_OUTPUT_EXTENSION}) - - # Initialize list of includes with those specified by the user. Append with - # ones specified to cmake directly. - set(HIP_HIPCC_INCLUDE_ARGS ${HIP_HIPCC_INCLUDE_ARGS_USER}) - - # Add the include directories - set(include_directories_generator "$") - list(APPEND HIP_HIPCC_INCLUDE_ARGS "$<$:-I$>") - - get_directory_property(_hip_include_directories INCLUDE_DIRECTORIES) - list(REMOVE_DUPLICATES _hip_include_directories) - if(_hip_include_directories) - foreach(dir ${_hip_include_directories}) - list(APPEND HIP_HIPCC_INCLUDE_ARGS $<$:-I${dir}>) - endforeach() - endif() - - HIP_GET_SOURCES_AND_OPTIONS(_hip_sources _hip_cmake_options _hipcc_options _hcc_options _nvcc_options ${ARGN}) - HIP_PARSE_HIPCC_OPTIONS(HIP_HIPCC_FLAGS ${_hipcc_options}) - HIP_PARSE_HIPCC_OPTIONS(HIP_HCC_FLAGS ${_hcc_options}) - HIP_PARSE_HIPCC_OPTIONS(HIP_NVCC_FLAGS ${_nvcc_options}) - - # Add the compile definitions - set(compile_definition_generator "$") - list(APPEND HIP_HIPCC_FLAGS "$<$:-D$>") - - # Check if we are building shared library. - set(_hip_build_shared_libs FALSE) - list(FIND _hip_cmake_options SHARED _hip_found_SHARED) - list(FIND _hip_cmake_options MODULE _hip_found_MODULE) - if(_hip_found_SHARED GREATER -1 OR _hip_found_MODULE GREATER -1) - set(_hip_build_shared_libs TRUE) - endif() - list(FIND _hip_cmake_options STATIC _hip_found_STATIC) - if(_hip_found_STATIC GREATER -1) - set(_hip_build_shared_libs FALSE) - endif() - - # If we are building a shared library, add extra flags to HIP_HIPCC_FLAGS - if(_hip_build_shared_libs) - list(APPEND HIP_HCC_FLAGS "-fPIC") - list(APPEND HIP_NVCC_FLAGS "--shared -Xcompiler '-fPIC'") - endif() - - # Set host compiler - set(HIP_HOST_COMPILER "${CMAKE_${HIP_C_OR_CXX}_COMPILER}") - - # Set compiler flags - set(_HIP_HOST_FLAGS "set(CMAKE_HOST_FLAGS ${CMAKE_${HIP_C_OR_CXX}_FLAGS})") - set(_HIP_HIPCC_FLAGS "set(HIP_HIPCC_FLAGS ${HIP_HIPCC_FLAGS})") - set(_HIP_HCC_FLAGS "set(HIP_HCC_FLAGS ${HIP_HCC_FLAGS})") - set(_HIP_NVCC_FLAGS "set(HIP_NVCC_FLAGS ${HIP_NVCC_FLAGS})") - foreach(config ${_hip_configuration_types}) - string(TOUPPER ${config} config_upper) - set(_HIP_HOST_FLAGS "${_HIP_HOST_FLAGS}\nset(CMAKE_HOST_FLAGS_${config_upper} ${CMAKE_${HIP_C_OR_CXX}_FLAGS_${config_upper}})") - set(_HIP_HIPCC_FLAGS "${_HIP_HIPCC_FLAGS}\nset(HIP_HIPCC_FLAGS_${config_upper} ${HIP_HIPCC_FLAGS_${config_upper}})") - set(_HIP_HCC_FLAGS "${_HIP_HCC_FLAGS}\nset(HIP_HCC_FLAGS_${config_upper} ${HIP_HCC_FLAGS_${config_upper}})") - set(_HIP_NVCC_FLAGS "${_HIP_NVCC_FLAGS}\nset(HIP_NVCC_FLAGS_${config_upper} ${HIP_NVCC_FLAGS_${config_upper}})") - endforeach() - - # Reset the output variable - set(_hip_generated_files "") - set(_hip_source_files "") - - # Iterate over all arguments and create custom commands for all source files - foreach(file ${ARGN}) - # Ignore any file marked as a HEADER_FILE_ONLY - get_source_file_property(_is_header ${file} HEADER_FILE_ONLY) - # Allow per source file overrides of the format. Also allows compiling non .cu files. - get_source_file_property(_hip_source_format ${file} HIP_SOURCE_PROPERTY_FORMAT) - if((${file} MATCHES "\\.cu$" OR _hip_source_format) AND NOT _is_header) - set(host_flag FALSE) - else() - set(host_flag TRUE) - endif() - - if(NOT host_flag) - # Determine output directory - HIP_COMPUTE_BUILD_PATH("${file}" hip_build_path) - set(hip_compile_output_dir "${CMAKE_CURRENT_BINARY_DIR}/CMakeFiles/${_target}.dir/${hip_build_path}") - - get_filename_component(basename ${file} NAME) - set(generated_file_path "${hip_compile_output_dir}/${CMAKE_CFG_INTDIR}") - set(generated_file_basename "${_target}_generated_${basename}${generated_extension}") - - # Set file names - set(generated_file "${generated_file_path}/${generated_file_basename}") - set(cmake_dependency_file "${hip_compile_output_dir}/${generated_file_basename}.depend") - set(custom_target_script_pregen "${hip_compile_output_dir}/${generated_file_basename}.cmake.pre-gen") - set(custom_target_script "${hip_compile_output_dir}/${generated_file_basename}.cmake") - - # Set properties for object files - set_source_files_properties("${generated_file}" - PROPERTIES - EXTERNAL_OBJECT true # This is an object file not to be compiled, but only be linked - ) - - # Don't add CMAKE_CURRENT_SOURCE_DIR if the path is already an absolute path - get_filename_component(file_path "${file}" PATH) - if(IS_ABSOLUTE "${file_path}") - set(source_file "${file}") - else() - set(source_file "${CMAKE_CURRENT_SOURCE_DIR}/${file}") - endif() - - # Bring in the dependencies - HIP_INCLUDE_HIPCC_DEPENDENCIES(${cmake_dependency_file}) - - # Configure the build script - configure_file("${HIP_run_hipcc}" "${custom_target_script_pregen}" @ONLY) - file(GENERATE - OUTPUT "${custom_target_script}" - INPUT "${custom_target_script_pregen}" - ) - set(main_dep DEPENDS ${source_file}) - if(CMAKE_GENERATOR MATCHES "Makefiles") - set(verbose_output "$(VERBOSE)") - elseif(HIP_VERBOSE_BUILD) - set(verbose_output ON) - else() - set(verbose_output OFF) - endif() - - # Create up the comment string - file(RELATIVE_PATH generated_file_relative_path "${CMAKE_BINARY_DIR}" "${generated_file}") - set(hip_build_comment_string "Building HIPCC object ${generated_file_relative_path}") - - # Build the generated file and dependency file - add_custom_command( - OUTPUT ${generated_file} - # These output files depend on the source_file and the contents of cmake_dependency_file - ${main_dep} - DEPENDS ${HIP_HIPCC_DEPEND} - DEPENDS ${custom_target_script} - # Make sure the output directory exists before trying to write to it. - COMMAND ${CMAKE_COMMAND} -E make_directory "${generated_file_path}" - COMMAND ${CMAKE_COMMAND} ARGS - -D verbose:BOOL=${verbose_output} - -D build_configuration:STRING=${_hip_build_configuration} - -D "generated_file:STRING=${generated_file}" - -P "${custom_target_script}" - WORKING_DIRECTORY "${hip_compile_output_dir}" - COMMENT "${hip_build_comment_string}" - ) - - # Make sure the build system knows the file is generated - set_source_files_properties(${generated_file} PROPERTIES GENERATED TRUE) - list(APPEND _hip_generated_files ${generated_file}) - list(APPEND _hip_source_files ${file}) - endif() - endforeach() - - # Set the return parameter - set(${_generated_files} ${_hip_generated_files}) - set(${_source_files} ${_hip_source_files}) -endmacro() - -############################################################################### -# HIP_ADD_EXECUTABLE -############################################################################### -macro(HIP_ADD_EXECUTABLE hip_target) - # Separate the sources from the options - HIP_GET_SOURCES_AND_OPTIONS(_sources _cmake_options _hipcc_options _hcc_options _nvcc_options ${ARGN}) - HIP_PREPARE_TARGET_COMMANDS(${hip_target} OBJ _generated_files _source_files ${_sources} HIPCC_OPTIONS ${_hipcc_options} HCC_OPTIONS ${_hcc_options} NVCC_OPTIONS ${_nvcc_options}) - if(_source_files) - list(REMOVE_ITEM _sources ${_source_files}) - endif() - if("x${HCC_HOME}" STREQUAL "x") - set(HCC_HOME "/opt/rocm/hcc") - endif() - set(CMAKE_HIP_LINK_EXECUTABLE "${HIP_HIPCC_CMAKE_LINKER_HELPER} ${HCC_HOME} -o ") - add_executable(${hip_target} ${_cmake_options} ${_generated_files} ${_sources}) - set_target_properties(${hip_target} PROPERTIES LINKER_LANGUAGE HIP) -endmacro() - -############################################################################### -# HIP_ADD_LIBRARY -############################################################################### -macro(HIP_ADD_LIBRARY hip_target) - # Separate the sources from the options - HIP_GET_SOURCES_AND_OPTIONS(_sources _cmake_options _hipcc_options _hcc_options _nvcc_options ${ARGN}) - HIP_PREPARE_TARGET_COMMANDS(${hip_target} OBJ _generated_files _source_files ${_sources} ${_cmake_options} HIPCC_OPTIONS ${_hipcc_options} HCC_OPTIONS ${_hcc_options} NVCC_OPTIONS ${_nvcc_options}) - if(_source_files) - list(REMOVE_ITEM _sources ${_source_files}) - endif() - add_library(${hip_target} ${_cmake_options} ${_generated_files} ${_sources}) - set_target_properties(${hip_target} PROPERTIES LINKER_LANGUAGE ${HIP_C_OR_CXX}) -endmacro() - -# vim: ts=4:sw=4:expandtab:smartindent diff --git a/CMake/HIP/FindHIP/run_hipcc.cmake b/CMake/HIP/FindHIP/run_hipcc.cmake deleted file mode 100644 index 4dc2572e98..0000000000 --- a/CMake/HIP/FindHIP/run_hipcc.cmake +++ /dev/null @@ -1,168 +0,0 @@ -############################################################################### -# Runs commands using HIPCC -############################################################################### - -############################################################################### -# This file runs the hipcc commands to produce the desired output file -# along with the dependency file needed by CMake to compute dependencies. -# -# Input variables: -# -# verbose:BOOL=<> OFF: Be as quiet as possible (default) -# ON : Describe each step -# build_configuration:STRING=<> Build configuration. Defaults to Debug. -# generated_file:STRING=<> File to generate. Mandatory argument. - -if(NOT build_configuration) - set(build_configuration Debug) -endif() -if(NOT generated_file) - message(FATAL_ERROR "You must specify generated_file on the command line") -endif() - -# Set these up as variables to make reading the generated file easier -set(HIP_HIPCC_EXECUTABLE "@HIP_HIPCC_EXECUTABLE@") # path -set(HIP_HIPCONFIG_EXECUTABLE "@HIP_HIPCONFIG_EXECUTABLE@") #path -set(HIP_HOST_COMPILER "@HIP_HOST_COMPILER@") # path -set(CMAKE_COMMAND "@CMAKE_COMMAND@") # path -set(HIP_run_make2cmake "@HIP_run_make2cmake@") # path -set(HCC_HOME "@HCC_HOME@") #path - -@HIP_HOST_FLAGS@ -@_HIP_HIPCC_FLAGS@ -@_HIP_HCC_FLAGS@ -@_HIP_NVCC_FLAGS@ -set(HIP_HIPCC_INCLUDE_ARGS "@HIP_HIPCC_INCLUDE_ARGS@") # list (needs to be in quotes to handle spaces properly) - -set(cmake_dependency_file "@cmake_dependency_file@") # path -set(source_file "@source_file@") # path -set(host_flag "@host_flag@") # bool - -# Determine compiler and compiler flags -execute_process(COMMAND ${HIP_HIPCONFIG_EXECUTABLE} --platform OUTPUT_VARIABLE HIP_PLATFORM OUTPUT_STRIP_TRAILING_WHITESPACE) -if(NOT host_flag) - set(__CC ${HIP_HIPCC_EXECUTABLE}) - if(HIP_PLATFORM STREQUAL "hcc") - if(NOT "x${HCC_HOME}" STREQUAL "x") - set(ENV{HCC_HOME} ${HCC_HOME}) - endif() - set(__CC_FLAGS ${HIP_HIPCC_FLAGS} ${HIP_HCC_FLAGS} ${HIP_HIPCC_FLAGS_${build_configuration}} ${HIP_HCC_FLAGS_${build_configuration}}) - else() - set(__CC_FLAGS ${HIP_HIPCC_FLAGS} ${HIP_NVCC_FLAGS} ${HIP_HIPCC_FLAGS_${build_configuration}} ${HIP_NVCC_FLAGS_${build_configuration}}) - endif() -else() - set(__CC ${HIP_HOST_COMPILER}) - set(__CC_FLAGS ${CMAKE_HOST_FLAGS} ${CMAKE_HOST_FLAGS_${build_configuration}}) -endif() -set(__CC_INCLUDES ${HIP_HIPCC_INCLUDE_ARGS}) - -# hip_execute_process - Executes a command with optional command echo and status message. -# status - Status message to print if verbose is true -# command - COMMAND argument from the usual execute_process argument structure -# ARGN - Remaining arguments are the command with arguments -# HIP_result - Return value from running the command -macro(hip_execute_process status command) - set(_command ${command}) - if(NOT "x${_command}" STREQUAL "xCOMMAND") - message(FATAL_ERROR "Malformed call to hip_execute_process. Missing COMMAND as second argument. (command = ${command})") - endif() - if(verbose) - execute_process(COMMAND "${CMAKE_COMMAND}" -E echo -- ${status}) - # Build command string to print - set(hip_execute_process_string) - foreach(arg ${ARGN}) - # Escape quotes if any - string(REPLACE "\"" "\\\"" arg ${arg}) - # Surround args with spaces with quotes - if(arg MATCHES " ") - list(APPEND hip_execute_process_string "\"${arg}\"") - else() - list(APPEND hip_execute_process_string ${arg}) - endif() - endforeach() - # Echo the command - execute_process(COMMAND ${CMAKE_COMMAND} -E echo ${hip_execute_process_string}) - endif() - # Run the command - execute_process(COMMAND ${ARGN} RESULT_VARIABLE HIP_result) -endmacro() - -# Delete the target file -hip_execute_process( - "Removing ${generated_file}" - COMMAND "${CMAKE_COMMAND}" -E remove "${generated_file}" - ) - -# Generate the dependency file -hip_execute_process( - "Generating dependency file: ${cmake_dependency_file}.pre" - COMMAND "${__CC}" - -M - "${source_file}" - -o "${cmake_dependency_file}.pre" - ${__CC_FLAGS} - ${__CC_INCLUDES} - ) - -if(HIP_result) - message(FATAL_ERROR "Error generating ${generated_file}") -endif() - -# Generate the cmake readable dependency file to a temp file -hip_execute_process( - "Generating temporary cmake readable file: ${cmake_dependency_file}.tmp" - COMMAND "${CMAKE_COMMAND}" - -D "input_file:FILEPATH=${cmake_dependency_file}.pre" - -D "output_file:FILEPATH=${cmake_dependency_file}.tmp" - -D "verbose=${verbose}" - -P "${HIP_run_make2cmake}" - ) - -if(HIP_result) - message(FATAL_ERROR "Error generating ${generated_file}") -endif() - -# Copy the file if it is different -hip_execute_process( - "Copy if different ${cmake_dependency_file}.tmp to ${cmake_dependency_file}" - COMMAND "${CMAKE_COMMAND}" -E copy_if_different "${cmake_dependency_file}.tmp" "${cmake_dependency_file}" - ) - -if(HIP_result) - message(FATAL_ERROR "Error generating ${generated_file}") -endif() - -# Delete the temporary file -hip_execute_process( - "Removing ${cmake_dependency_file}.tmp and ${cmake_dependency_file}.pre" - COMMAND "${CMAKE_COMMAND}" -E remove "${cmake_dependency_file}.tmp" "${cmake_dependency_file}.pre" - ) - -if(HIP_result) - message(FATAL_ERROR "Error generating ${generated_file}") -endif() - -# Generate the output file -hip_execute_process( - "Generating ${generated_file}" - COMMAND "${__CC}" - -c - "${source_file}" - -o "${generated_file}" - ${__CC_FLAGS} - ${__CC_INCLUDES} - ) - -if(HIP_result) - # Make sure that we delete the output file - hip_execute_process( - "Removing ${generated_file}" - COMMAND "${CMAKE_COMMAND}" -E remove "${generated_file}" - ) - message(FATAL_ERROR "Error generating file ${generated_file}") -else() - if(verbose) - message("Generated ${generated_file} successfully.") - endif() -endif() -# vim: ts=4:sw=4:expandtab:smartindent diff --git a/CMake/HIP/FindHIP/run_make2cmake.cmake b/CMake/HIP/FindHIP/run_make2cmake.cmake deleted file mode 100644 index d2e3eb5169..0000000000 --- a/CMake/HIP/FindHIP/run_make2cmake.cmake +++ /dev/null @@ -1,50 +0,0 @@ -############################################################################### -# Computes dependencies using HIPCC -############################################################################### - -############################################################################### -# This file converts dependency files generated using hipcc to a format that -# cmake can understand. - -# Input variables: -# -# input_file:STRING=<> Dependency file to parse. Required argument -# output_file:STRING=<> Output file to generate. Required argument - -if(NOT input_file OR NOT output_file) - message(FATAL_ERROR "You must specify input_file and output_file on the command line") -endif() - -file(READ ${input_file} depend_text) - -if (NOT "${depend_text}" STREQUAL "") - string(REPLACE " /" "\n/" depend_text ${depend_text}) - string(REGEX REPLACE "^.*:" "" depend_text ${depend_text}) - string(REGEX REPLACE "[ \\\\]*\n" ";" depend_text ${depend_text}) - - set(dependency_list "") - - foreach(file ${depend_text}) - string(REGEX REPLACE "^ +" "" file ${file}) - if(NOT EXISTS "${file}") - message(WARNING " Removing non-existent dependency file: ${file}") - set(file "") - endif() - - if(NOT IS_DIRECTORY "${file}") - get_filename_component(file_absolute "${file}" ABSOLUTE) - list(APPEND dependency_list "${file_absolute}") - endif() - endforeach() -endif() - -# Remove the duplicate entries and sort them. -list(REMOVE_DUPLICATES dependency_list) -list(SORT dependency_list) - -foreach(file ${dependency_list}) - set(hip_hipcc_depend "${hip_hipcc_depend} \"${file}\"\n") -endforeach() - -file(WRITE ${output_file} "# Generated by: FindHIP.cmake. Do not edit.\nSET(HIP_HIPCC_DEPEND\n ${hip_hipcc_depend})\n\n") -# vim: ts=4:sw=4:expandtab:smartindent diff --git a/CMake/hoomd/FindCUDALibs.cmake b/CMake/hoomd/FindCUDALibs.cmake index 65d183011c..deb96ab125 100644 --- a/CMake/hoomd/FindCUDALibs.cmake +++ b/CMake/hoomd/FindCUDALibs.cmake @@ -1,11 +1,11 @@ # Find CUDA libraries and binaries used by HOOMD -# find CUDA library path -get_filename_component(CUDA_BIN_PATH ${CMAKE_CUDA_COMPILER} DIRECTORY) -get_filename_component(CUDA_LIB_PATH "${CUDA_BIN_PATH}/../lib64/" ABSOLUTE) - set(REQUIRED_CUDA_LIB_VARS "") if (HIP_PLATFORM STREQUAL "nvcc") + # find CUDA library path + get_filename_component(CUDA_BIN_PATH ${CMAKE_CUDA_COMPILER} DIRECTORY) + get_filename_component(CUDA_LIB_PATH "${CUDA_BIN_PATH}/../lib64/" ABSOLUTE) + # find libraries that go with this compiler find_library(CUDA_cudart_LIBRARY cudart HINTS ${CUDA_LIB_PATH}) mark_as_advanced(CUDA_cudart_LIBRARY) @@ -22,42 +22,6 @@ else() add_library(CUDA::cudart UNKNOWN IMPORTED) endif() -if (HIP_PLATFORM STREQUAL "hip-clang" OR HIP_PLATFORM STREQUAL "hcc") - # find libraries that go with this compiler - find_library(HIP_hip_hcc_LIBRARY hip_hcc - PATHS - "${HIP_ROOT_DIR}" - ENV ROCM_PATH - ENV HIP_PATH - /opt/rocm - PATH_SUFFIXES lib - NO_DEFAULT_PATH) - mark_as_advanced(HIP_hip_hcc_LIBRARY) - find_library(HIP_hiprtc_LIBRARY hiprtc - PATHS - "${HIP_ROOT_DIR}" - ENV ROCM_PATH - ENV HIP_PATH - /opt/rocm - PATH_SUFFIXES lib - NO_DEFAULT_PATH) - mark_as_advanced(HIP_hiprtc_LIBRARY) - - if(HIP_hip_hcc_LIBRARY AND NOT TARGET HIP::hiprt) - add_library(HIP::hiprt UNKNOWN IMPORTED) - set_target_properties(HIP::hiprt PROPERTIES - IMPORTED_LOCATION "${HIP_hip_hcc_LIBRARY}" - INTERFACE_LINK_LIBRARIES ${HIP_hiprtc_LIBRARY} - ) - endif() - list(APPEND REQUIRED_HIP_LIB_VARS "HIP_hip_hcc_LIBRARY") - list(APPEND REQUIRED_HIP_LIB_VARS "HIP_hiprtc_LIBRARY") -else() - # define empty target - add_library(HIP::hiprt UNKNOWN IMPORTED) -endif() - - if (HIP_PLATFORM STREQUAL "nvcc") find_library(CUDA_cudadevrt_LIBRARY cudadevrt HINTS ${CUDA_LIB_PATH}) mark_as_advanced(CUDA_cudadevrt_LIBRARY) @@ -197,83 +161,6 @@ else() add_library(CUDA::cusparse UNKNOWN IMPORTED) endif() -if (HIP_PLATFORM STREQUAL "hip-clang" OR HIP_PLATFORM STREQUAL "hcc") - find_path(HIP_hipfft_INCLUDE_DIR - NAMES hipfft.h - PATHS - ${HIP_ROOT_DIR}/rocfft/include - $ENV{ROCM_PATH}/hipfft/include - $ENV{HIP_PATH}/hipfft/include - /opt/rocm/include - /opt/rocm/hipfft/include - NO_DEFAULT_PATH) - - list(APPEND REQUIRED_CUDA_LIB_VARS HIP_hipfft_INCLUDE_DIR) - - find_library(HIP_rocfft_LIBRARY rocfft - PATHS - "${HIP_ROOT_DIR}" - $ENV{ROCM_PATH}/rocfft - ENV HIP_PATH - /opt/rocm - /opt/rocm/rocfft - PATH_SUFFIXES lib - NO_DEFAULT_PATH) - - find_path(HIP_rocfft_INCLUDE_DIR - NAMES rocfft.h - PATHS - ${HIP_ROOT_DIR}/rocfft - $ENV{ROCM_PATH}/rocfft - $ENV{HIP_PATH}/rocfft - /opt/rocm - /opt/rocm/rocfft - PATH_SUFFIXES include - NO_DEFAULT_PATH) - - mark_as_advanced(HIP_rocfft_LIBRARY) - if(HIP_rocfft_LIBRARY AND NOT TARGET HIP::hipfft) - add_library(HIP::hipfft UNKNOWN IMPORTED) - set_target_properties(HIP::hipfft PROPERTIES - IMPORTED_LOCATION "${HIP_rocfft_LIBRARY}" - INTERFACE_INCLUDE_DIRECTORIES "${HIP_hipfft_INCLUDE_DIR};${HIP_rocfft_INCLUDE_DIR}" - ) - endif() - list(APPEND REQUIRED_CUDA_LIB_VARS HIP_rocfft_LIBRARY) -endif() - -if (HIP_PLATFORM STREQUAL "hip-clang" OR HIP_PLATFORM STREQUAL "hcc") - find_library(HIP_roctracer_LIBRARY roctracer64 - PATHS - "${HIP_ROOT_DIR}" - ENV ROCM_PATH - ENV HIP_PATH - /opt/rocm - /opt/rocm/roctracer - PATH_SUFFIXES lib - NO_DEFAULT_PATH) - - find_path(HIP_roctracer_INCLUDE_DIR - NAMES roctracer.h - PATHS - ${HIP_ROOT_DIR}/roctracer - $ENV{ROCM_PATH}/roctracer - $ENV{HIP_PATH}/roctracer - /opt/rocm - /opt/rocm/roctracer - PATH_SUFFIXES include - NO_DEFAULT_PATH) - - mark_as_advanced(HIP_roctracer_LIBRARY) - if(HIP_roctracer_LIBRARY AND NOT TARGET HIP::roctracer) - add_library(HIP::roctracer UNKNOWN IMPORTED) - set_target_properties(HIP::roctracer PROPERTIES - IMPORTED_LOCATION "${HIP_roctracer_LIBRARY}" - INTERFACE_INCLUDE_DIRECTORIES "${HIP_roctracer_INCLUDE_DIR};${HIP_roctracer_INCLUDE_DIR}" - ) - endif() -endif() - #find_library(HIP_hipsparse_LIBRARY hipsparse # PATHS @@ -304,7 +191,8 @@ endif() # ) #endif() -#if (HIP_PLATFORM STREQUAL "hip-clang" OR HIP_PLATFORM STREQUAL "hcc") +if (HIP_PLATFORM STREQUAL "hip-clang") + find_package(hipfft) # find_library(HIP_rocsparse_LIBRARY rocsparse # PATHS # "${HIP_ROOT_DIR}" @@ -331,7 +219,7 @@ endif() # INTERFACE_LINK_LIBRARIES "${HIP_rocsparse_LIBRARY}" # ) # list(APPEND REQUIRED_CUDA_LIB_VARS HIP_rocsparse_LIBRARY) -#endif() +endif() if (HIP_PLATFORM STREQUAL "nvcc") # find cuda-memcheck diff --git a/CMake/hoomd/HOOMDHIPSetup.cmake b/CMake/hoomd/HOOMDHIPSetup.cmake index 2ea2873c73..9c14f9a64f 100644 --- a/CMake/hoomd/HOOMDHIPSetup.cmake +++ b/CMake/hoomd/HOOMDHIPSetup.cmake @@ -1,92 +1,44 @@ if(ENABLE_HIP) - find_package(HIP QUIET) + find_package(HIP) if (HIP_FOUND) - # call hipcc to tell us about the backend compiler - set(ENV{HIPCC_VERBOSE} 1) - - FILE(WRITE ${CMAKE_CURRENT_BINARY_DIR}/hip_test.cc " -int main(int argc, char **argv) -{ } -") - EXECUTE_PROCESS(COMMAND ${HIP_HIPCC_EXECUTABLE} ${CMAKE_CURRENT_BINARY_DIR}/hip_test.cc OUTPUT_VARIABLE _hipcc_verbose_out) - - string(REPLACE " " ";" _hipcc_verbose_options ${_hipcc_verbose_out}) - - # get the compiler executable for device code - LIST(GET _hipcc_verbose_options 1 _hip_compiler) - - # set it as the compiler - if (${_hip_compiler} MATCHES nvcc) - set(HIP_PLATFORM nvcc) - elseif(${_hip_compiler} MATCHES hcc) - set(HIP_PLATFORM hcc) - elseif(${_hip_compiler} MATCHES clang) - set(HIP_PLATFORM hip-clang) - else() - message(ERROR "Unknown HIP backend " ${_hip_compiler}) - endif() - - # use hipcc as C++ linker for shared libraries - SET(CMAKE_CUDA_COMPILER ${HIP_HIPCC_EXECUTABLE}) - - # this is hack to set the right options on hipcc, may not be portable - include(hipcc) - - # override command line, so that it doesn't contain "-x cu" - set(CMAKE_CUDA_COMPILE_WHOLE_COMPILATION - " ${CMAKE_CUDA_HOST_FLAGS} -c -o ") + ENABLE_LANGUAGE(HIP) # setup nvcc to build for all CUDA architectures. Allow user to modify the list if desired - set(AMDGPU_TARGET_LIST gfx900 gfx906 gfx908 CACHE STRING "List of AMD GPU to compile HIP code for. Separate with semicolons.") - - foreach(_amdgpu_target ${AMDGPU_TARGET_LIST}) - set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} --amdgpu-target=${_amdgpu_target}") - endforeach (_amdgpu_target) - - if (HIP_FOUND) - # reduce link time (no device linking) - set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} -fno-gpu-rdc") - endif() + set(CMAKE_HIP_ARCHITECTURES gfx900 gfx906 gfx908 gfx90a CACHE STRING "List of AMD GPU to compile HIP code for. Separate with semicolons.") + set(HIP_PLATFORM hip-clang) - # these are no-ops, as device linking is not supported with hcc - set(CMAKE_CUDA_DEVICE_LINK_LIBRARY " -o -x c++ -c /dev/null") - set(CMAKE_CUDA_DEVICE_LINK_EXECUTABLE " -o -x c++ -c /dev/null") + #foreach(_amdgpu_target ${AMDGPU_TARGET_LIST}) + # set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} --amdgpu-target=${_amdgpu_target}") + #endforeach (_amdgpu_target) - if(CMAKE_GENERATOR STREQUAL "Ninja") - # this is also ugly, but ninja/hipcc is only supported with a future cmake - CMAKE_MINIMUM_REQUIRED(VERSION 3.17.0 FATAL_ERROR) - - # hipcc can write dependencies (undocumented CMake option) - set(CMAKE_DEPFILE_FLAGS_CUDA "-MD -MT -MF ") - endif() - - # don't let CMake examine the compiler, because it will fail - SET(CMAKE_CUDA_COMPILER_FORCED TRUE) + #if (HIP_FOUND) + # # reduce link time (no device linking) + # set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} -fno-gpu-rdc") + #endif() #search for HIP include directory - find_path(HIP_INCLUDE_DIR hip/hip_runtime.h - PATHS - "${HIP_ROOT_DIR}" - ENV ROCM_PATH - ENV HIP_PATH - PATH_SUFFIXES include) - - find_path(ROCm_hsa_INCLUDE_DIR - NAMES hsa/hsa.h - PATHS - ${HIP_ROOT_DIR}/hsa - ${HIP_ROOT_DIR}/hsa - $ENV{ROCM_PATH}/hsa - $ENV{HIP_PATH}/hsa - $ENV{HSA_PATH} - /opt/rocm - PATH_SUFFIXES include - NO_DEFAULT_PATH) - - option(ENABLE_ROCTRACER "Enable roctracer profiler integration" off) - - list(APPEND HIP_INCLUDE_DIR ${ROCm_hsa_INCLUDE_DIR}) +# find_path(HIP_INCLUDE_DIR hip/hip_runtime.h +# PATHS +# "${HIP_ROOT_DIR}" +# ENV ROCM_PATH +# ENV HIP_PATH +# PATH_SUFFIXES include) +# +# find_path(ROCm_hsa_INCLUDE_DIR +# NAMES hsa/hsa.h +# PATHS +# ${HIP_ROOT_DIR}/hsa +# ${HIP_ROOT_DIR}/hsa +# $ENV{ROCM_PATH}/hsa +# $ENV{HIP_PATH}/hsa +# $ENV{HSA_PATH} +# /opt/rocm +# PATH_SUFFIXES include +# NO_DEFAULT_PATH) +# +# +# list(APPEND HIP_INCLUDE_DIR ${ROCm_hsa_INCLUDE_DIR}) else() # here we go if hipcc is not available, fall back on internal HIP->CUDA headers ENABLE_LANGUAGE(CUDA) @@ -112,29 +64,35 @@ int main(int argc, char **argv) endif() endif() - ENABLE_LANGUAGE(CUDA) - - if(NOT TARGET HIP::hip) - add_library(HIP::hip INTERFACE IMPORTED) - set_target_properties(HIP::hip PROPERTIES - INTERFACE_INCLUDE_DIRECTORIES "${HIP_INCLUDE_DIR};${HIPCUB_INCLUDE_DIR}") - - # set HIP_VERSION_* on non-CUDA targets (the version is already defined on AMD targets through hipcc) - set_property(TARGET HIP::hip APPEND PROPERTY INTERFACE_COMPILE_DEFINITIONS - $<$>:HIP_VERSION_MAJOR=${HIP_VERSION_MAJOR}>) - set_property(TARGET HIP::hip APPEND PROPERTY INTERFACE_COMPILE_DEFINITIONS - $<$>:HIP_VERSION_MINOR=${HIP_VERSION_MINOR}>) - set_property(TARGET HIP::hip APPEND PROPERTY INTERFACE_COMPILE_DEFINITIONS - $<$>:HIP_VERSION_PATCH=${HIP_VERSION_PATCH}>) - - # branch upon HCC or NVCC target - if(${HIP_PLATFORM} STREQUAL "nvcc") - set_property(TARGET HIP::hip APPEND PROPERTY INTERFACE_COMPILE_DEFINITIONS - $<$>:__HIP_PLATFORM_NVCC__>) - elseif(${HIP_PLATFORM} STREQUAL "hcc" OR ${HIP_PLATFORM} STREQUAL "hip-clang") - set_property(TARGET HIP::hip APPEND PROPERTY INTERFACE_COMPILE_DEFINITIONS - $<$>:__HIP_PLATFORM_HCC__>) - endif() +# if(NOT TARGET HIP::hip) +# add_library(HIP::hip INTERFACE IMPORTED) +# set_target_properties(HIP::hip PROPERTIES +# INTERFACE_INCLUDE_DIRECTORIES "${HIP_INCLUDE_DIR};${HIPCUB_INCLUDE_DIR}") +# +# # set HIP_VERSION_* on non-CUDA targets (the version is already defined on AMD targets through hipcc) +# set_property(TARGET HIP::hip APPEND PROPERTY INTERFACE_COMPILE_DEFINITIONS +# $<$>:HIP_VERSION_MAJOR=${HIP_VERSION_MAJOR}>) +# set_property(TARGET HIP::hip APPEND PROPERTY INTERFACE_COMPILE_DEFINITIONS +# $<$>:HIP_VERSION_MINOR=${HIP_VERSION_MINOR}>) +# set_property(TARGET HIP::hip APPEND PROPERTY INTERFACE_COMPILE_DEFINITIONS +# $<$>:HIP_VERSION_PATCH=${HIP_VERSION_PATCH}>) +# +# # branch upon HCC or NVCC target +# if(${HIP_PLATFORM} STREQUAL "nvcc") +# set_property(TARGET HIP::hip APPEND PROPERTY INTERFACE_COMPILE_DEFINITIONS +# $<$>:__HIP_PLATFORM_NVCC__>) +# elseif(${HIP_PLATFORM} STREQUAL "hip-clang") +# set_property(TARGET HIP::hip APPEND PROPERTY INTERFACE_COMPILE_DEFINITIONS +# $<$>:__HIP_PLATFORM_HCC__>) +# endif() +# endif() + + # branch upon HCC or NVCC target + if(${HIP_PLATFORM} STREQUAL "nvcc") + set_property(TARGET hip::host APPEND PROPERTY INTERFACE_COMPILE_DEFINITIONS __HIP_PLATFORM_NVCC__) + elseif(${HIP_PLATFORM} STREQUAL "hip-clang") + set_property(TARGET hip::host APPEND PROPERTY INTERFACE_COMPILE_DEFINITIONS __HIP_PLATFORM_AMD__) endif() + find_package(CUDALibs REQUIRED) endif() diff --git a/CMake/hoomd/hipcc.cmake b/CMake/hoomd/hipcc.cmake deleted file mode 100644 index 08907f7913..0000000000 --- a/CMake/hoomd/hipcc.cmake +++ /dev/null @@ -1,51 +0,0 @@ -set(CMAKE_CUDA_COMPILER_HAS_DEVICE_LINK_PHASE True) -set(CMAKE_CUDA_VERBOSE_FLAG "-v") -set(CMAKE_CUDA_VERBOSE_COMPILE_FLAG "-v") - -if(NOT "x${CMAKE_CUDA_SIMULATE_ID}" STREQUAL "xMSVC") - set(CMAKE_CUDA_COMPILE_OPTIONS_PIE -fPIE) - set(CMAKE_CUDA_COMPILE_OPTIONS_PIC -fPIC) - - # hipcc targets should have default visibility - set(CMAKE_CUDA_COMPILE_OPTIONS_VISIBILITY -fvisibility=) - # CMAKE_SHARED_LIBRARY_CUDA_FLAGS is sent to the host linker so we - # don't need to forward it through nvcc. - set(CMAKE_SHARED_LIBRARY_CUDA_FLAGS -fPIC) - string(APPEND CMAKE_CUDA_FLAGS_INIT " ") - string(APPEND CMAKE_CUDA_FLAGS_DEBUG_INIT " -g") - string(APPEND CMAKE_CUDA_FLAGS_RELEASE_INIT " -O3 -DNDEBUG") - string(APPEND CMAKE_CUDA_FLAGS_MINSIZEREL_INIT " -O1 -DNDEBUG") - string(APPEND CMAKE_CUDA_FLAGS_RELWITHDEBINFO_INIT " -O2 -g -DNDEBUG") -endif() -set(CMAKE_SHARED_LIBRARY_CREATE_CUDA_FLAGS -shared) -set(CMAKE_INCLUDE_SYSTEM_FLAG_CUDA -isystem) - -if("x${CMAKE_CUDA_SIMULATE_ID}" STREQUAL "xMSVC") - set(CMAKE_CUDA_STANDARD_DEFAULT "") -else() - set(CMAKE_CUDA_STANDARD_DEFAULT 11) - set(CMAKE_CUDA98_STANDARD_COMPILE_OPTION "") - set(CMAKE_CUDA98_EXTENSION_COMPILE_OPTION "") - - set(CMAKE_CUDA11_STANDARD_COMPILE_OPTION "-std=c++11") - set(CMAKE_CUDA11_EXTENSION_COMPILE_OPTION "-std=c++11") - - if (NOT CMAKE_CUDA_COMPILER_VERSION VERSION_LESS 9.0) - set(CMAKE_CUDA98_STANDARD_COMPILE_OPTION "-std=c++03") - set(CMAKE_CUDA98_EXTENSION_COMPILE_OPTION "-std=c++03") - set(CMAKE_CUDA14_STANDARD_COMPILE_OPTION "-std=c++14") - set(CMAKE_CUDA14_EXTENSION_COMPILE_OPTION "-std=c++14") - endif() - -endif() - -# FIXME: investigate use of --options-file. -# Tell Makefile generator that nvcc does not support @ syntax. -set(CMAKE_CUDA_USE_RESPONSE_FILE_FOR_INCLUDES 0) -set(CMAKE_CUDA_USE_RESPONSE_FILE_FOR_LIBRARIES 0) -set(CMAKE_CUDA_USE_RESPONSE_FILE_FOR_OBJECTS 0) - -if (CMAKE_CUDA_COMPILER_VERSION VERSION_GREATER_EQUAL "9.0") - set(CMAKE_CUDA_RESPONSE_FILE_LINK_FLAG "--options-file ") - set(CMAKE_CUDA_RESPONSE_FILE_FLAG "--options-file ") -endif() diff --git a/CMakeLists.txt b/CMakeLists.txt index df0ed9364f..812a174b61 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -62,6 +62,7 @@ endif() set(CMAKE_CXX_STANDARD 17) set(CMAKE_CXX_STANDARD_REQUIRED ON) set(CMAKE_CUDA_STANDARD 14) +set(CMAKE_HIP_STANDARD 14) # Enable compiler warnings on gcc and clang (common compilers used by developers) if(CMAKE_COMPILER_IS_GNUCXX OR CMAKE_CXX_COMPILER_ID MATCHES "Clang") @@ -207,13 +208,7 @@ configure_package_config_file(hoomd-config.cmake.in install(FILES CMake/hoomd/FindTBB.cmake CMake/hoomd/FindCUDALibs.cmake - CMake/HIP/FindHIP.cmake CMake/hoomd/HOOMDHIPSetup.cmake - CMake/hoomd/hipcc.cmake CMake/hoomd/hoomd-macros.cmake ${HOOMD_BINARY_DIR}/hoomd-config.cmake DESTINATION ${CONFIG_INSTALL_DIR}) - -install(FILES CMake/HIP/FindHIP/run_hipcc.cmake - CMake/HIP/FindHIP/run_make2cmake.cmake - DESTINATION ${CONFIG_INSTALL_DIR}/FindHIP) diff --git a/hoomd/CMakeLists.txt b/hoomd/CMakeLists.txt index d8c068b526..45f5b39351 100644 --- a/hoomd/CMakeLists.txt +++ b/hoomd/CMakeLists.txt @@ -268,8 +268,8 @@ endif(SINGLE_PRECISION) # Libraries and compile definitions for CUDA enabled builds if (ENABLE_HIP) - if (HIP_PLATFORM STREQUAL "hcc" OR HIP_PLATFORM STREQUAL "hip-clang") - target_link_libraries(_hoomd PUBLIC HIP::hipfft HIP::hiprt) + if (HIP_PLATFORM STREQUAL "hip-clang") + target_link_libraries(_hoomd PUBLIC hip::hipfft) elseif(HIP_PLATFORM STREQUAL "nvcc") target_link_libraries(_hoomd PUBLIC CUDA::cudart CUDA::cufft) endif() @@ -288,7 +288,7 @@ if (ENABLE_HIP) target_compile_definitions(_hoomd PUBLIC CUSOLVER_AVAILABLE) endif() - target_link_libraries(_hoomd PUBLIC HIP::hip) + target_link_libraries(_hoomd PUBLIC hip::host) if (ENABLE_ROCTRACER) target_link_libraries(_hoomd PUBLIC HIP::roctracer) diff --git a/hoomd/SnapshotSystemData.cc b/hoomd/SnapshotSystemData.cc index 26b0f0a5b6..2073f28077 100644 --- a/hoomd/SnapshotSystemData.cc +++ b/hoomd/SnapshotSystemData.cc @@ -52,7 +52,7 @@ template void SnapshotSystemData::wrap() auto const img = make_int3(static_cast(std::floor(frac.x)), static_cast(std::floor(frac.y)), static_cast(std::floor(frac.z))); - particle_data.image[i] += img; + particle_data.image[i] = particle_data.image[i] + img; } } diff --git a/hoomd/md/BondTablePotential.cc b/hoomd/md/BondTablePotential.cc index b7e08f61f5..6e99279143 100644 --- a/hoomd/md/BondTablePotential.cc +++ b/hoomd/md/BondTablePotential.cc @@ -128,8 +128,8 @@ pybind11::dict BondTablePotential::getParams(std::string type) auto type_id = m_bond_data->getTypeByName(type); pybind11::dict params; - params["r_min"] = h_params.data[type_id].x; - params["r_max"] = h_params.data[type_id].y; + params["r_min"] = (Scalar)h_params.data[type_id].x; + params["r_max"] = (Scalar)h_params.data[type_id].y; auto V = pybind11::array_t(m_table_width); auto V_unchecked = V.mutable_unchecked<1>(); diff --git a/hoomd/md/EvaluatorTersoff.h b/hoomd/md/EvaluatorTersoff.h index f415ce2e69..8611b3befd 100644 --- a/hoomd/md/EvaluatorTersoff.h +++ b/hoomd/md/EvaluatorTersoff.h @@ -91,15 +91,15 @@ class EvaluatorTersoff pybind11::dict v; pybind11::list mags; - mags.append(coeffs.x); - mags.append(coeffs.y); + mags.append((Scalar)coeffs.x); + mags.append((Scalar)coeffs.y); v["magnitudes"] = pybind11::tuple(mags); v["cutoff_thickness"] = cutoff_thickness; pybind11::list exp_factors; - exp_factors.append(exp_consts.x); - exp_factors.append(exp_consts.y); + exp_factors.append((Scalar)exp_consts.x); + exp_factors.append((Scalar)exp_consts.y); v["exp_factors"] = pybind11::tuple(exp_factors); v["lambda3"] = pow(lambda_cube, 1. / 3.); diff --git a/hoomd/md/ForceComposite.h b/hoomd/md/ForceComposite.h index ac39ba6a4f..7d633386ae 100644 --- a/hoomd/md/ForceComposite.h +++ b/hoomd/md/ForceComposite.h @@ -172,13 +172,13 @@ class PYBIND11_EXPORT ForceComposite : public MolecularForceCompute for (unsigned int i = 0; i < N; i++) { auto index = m_body_idx(body_type_id, i); - positions.append(pybind11::make_tuple(h_body_pos.data[index].x, - h_body_pos.data[index].y, - h_body_pos.data[index].z)); - orientations.append(pybind11::make_tuple(h_body_orientation.data[index].x, - h_body_orientation.data[index].y, - h_body_orientation.data[index].z, - h_body_orientation.data[index].w)); + positions.append(pybind11::make_tuple((Scalar)h_body_pos.data[index].x, + (Scalar)h_body_pos.data[index].y, + (Scalar)h_body_pos.data[index].z)); + orientations.append(pybind11::make_tuple((Scalar)h_body_orientation.data[index].x, + (Scalar)h_body_orientation.data[index].y, + (Scalar)h_body_orientation.data[index].z, + (Scalar)h_body_orientation.data[index].w)); types.append(m_pdata->getNameByType(h_body_types.data[index])); charges.append(m_body_charge[body_type_id][i]); diameters.append(m_body_diameter[body_type_id][i]); diff --git a/hoomd/md/OPLSDihedralForceCompute.cc b/hoomd/md/OPLSDihedralForceCompute.cc index 017f5626eb..bb366743e5 100644 --- a/hoomd/md/OPLSDihedralForceCompute.cc +++ b/hoomd/md/OPLSDihedralForceCompute.cc @@ -323,22 +323,22 @@ void OPLSDihedralForceCompute::computeForces(uint64_t timestep) f3.w = e_dihedral; // Apply force to each of the 4 atoms - h_force.data[i1].x += f1.x; - h_force.data[i1].y += f1.y; - h_force.data[i1].z += f1.z; - h_force.data[i1].w += f1.w; - h_force.data[i2].x += f2.x; - h_force.data[i2].y += f2.y; - h_force.data[i2].z += f2.z; - h_force.data[i2].w += f2.w; - h_force.data[i3].x += f3.x; - h_force.data[i3].y += f3.y; - h_force.data[i3].z += f3.z; - h_force.data[i3].w += f3.w; - h_force.data[i4].x += f4.x; - h_force.data[i4].y += f4.y; - h_force.data[i4].z += f4.z; - h_force.data[i4].w += f4.w; + h_force.data[i1].x = h_force.data[i1].x + f1.x; + h_force.data[i1].y = h_force.data[i1].y + f1.y; + h_force.data[i1].z = h_force.data[i1].z + f1.z; + h_force.data[i1].w = h_force.data[i1].w + f1.w; + h_force.data[i2].x = h_force.data[i2].x + f2.x; + h_force.data[i2].y = h_force.data[i2].y + f2.y; + h_force.data[i2].z = h_force.data[i2].z + f2.z; + h_force.data[i2].w = h_force.data[i2].w + f2.w; + h_force.data[i3].x = h_force.data[i3].x + f3.x; + h_force.data[i3].y = h_force.data[i3].y + f3.y; + h_force.data[i3].z = h_force.data[i3].z + f3.z; + h_force.data[i3].w = h_force.data[i3].w + f3.w; + h_force.data[i4].x = h_force.data[i4].x + f4.x; + h_force.data[i4].y = h_force.data[i4].y + f4.y; + h_force.data[i4].z = h_force.data[i4].z + f4.z; + h_force.data[i4].w = h_force.data[i4].z + f4.w; // Compute 1/4 of the virial, 1/4 for each atom in the dihedral // upper triangular version of virial tensor diff --git a/hoomd/md/TwoStepRATTLEBDGPU.h b/hoomd/md/TwoStepRATTLEBDGPU.h index fd90af349c..f9bc4a27a1 100644 --- a/hoomd/md/TwoStepRATTLEBDGPU.h +++ b/hoomd/md/TwoStepRATTLEBDGPU.h @@ -144,6 +144,7 @@ template void TwoStepRATTLEBDGPU::integrateStepOne(uin bool aniso = this->m_aniso; + #if defined(__HIP_PLATFORM_NVCC__) if (this->m_exec_conf->allConcurrentManagedAccess()) { // prefetch gammas @@ -160,6 +161,7 @@ template void TwoStepRATTLEBDGPU::integrateStepOne(uin if (this->m_exec_conf->isCUDAErrorCheckingEnabled()) CHECK_CUDA_ERROR(); } + #endif this->m_exec_conf->beginMultiGPU(); @@ -231,6 +233,7 @@ template void TwoStepRATTLEBDGPU::includeRATTLEForce(u args.timestep = timestep; args.seed = this->m_sysdef->getSeed(); + #if defined(__HIP_PLATFORM_NVCC__) if (this->m_exec_conf->allConcurrentManagedAccess()) { // prefetch gammas @@ -244,6 +247,7 @@ template void TwoStepRATTLEBDGPU::includeRATTLEForce(u if (this->m_exec_conf->isCUDAErrorCheckingEnabled()) CHECK_CUDA_ERROR(); } + #endif this->m_exec_conf->beginMultiGPU(); From 32eabb3928318058935830aaf26ce1c3bc78c73a Mon Sep 17 00:00:00 2001 From: Jens Glaser Date: Thu, 16 Jun 2022 09:55:29 -0400 Subject: [PATCH 02/15] compile HIP device code --- hoomd/CMakeLists.txt | 1 + hoomd/md/CMakeLists.txt | 11 +++++++++++ 2 files changed, 12 insertions(+) diff --git a/hoomd/CMakeLists.txt b/hoomd/CMakeLists.txt index 45f5b39351..a345309de9 100644 --- a/hoomd/CMakeLists.txt +++ b/hoomd/CMakeLists.txt @@ -212,6 +212,7 @@ set(_hoomd_cu_sources BondedGroupData.cu if (ENABLE_HIP) set(_cuda_sources ${_hoomd_cu_sources}) +set_source_files_properties(${_hoomd_cu_sources} PROPERTIES LANGUAGE HIP) endif (ENABLE_HIP) ######################### diff --git a/hoomd/md/CMakeLists.txt b/hoomd/md/CMakeLists.txt index 62fb919cce..3749b69ff9 100644 --- a/hoomd/md/CMakeLists.txt +++ b/hoomd/md/CMakeLists.txt @@ -274,6 +274,7 @@ set(_md_cu_sources ActiveForceComputeGPU.cu if (ENABLE_HIP) set(_cuda_sources ${_md_cu_sources} ${DFFT_CU_SOURCES}) +set_source_files_properties(${_md_cu_sources} PROPERTIES LANGUAGE HIP) endif (ENABLE_HIP) # generate pybind11 export cc files @@ -315,6 +316,7 @@ foreach(_manifold ${_manifolds}) ActiveForceConstraintCompute${_manifold}GPU.cu TwoStepRATTLE${_manifold}GPU.cu ) + set_source_files_properties(${_cuda_sources} PROPERTIES LANGUAGE HIP) endif() foreach (_method ${_rattle_methods}) @@ -328,6 +330,7 @@ foreach(_manifold ${_manifolds}) export_TwoStepRATTLE${_method}${_manifold}GPU.cc @ONLY) set(_cuda_sources ${_cuda_sources} export_TwoStepRATTLE${_method}${_manifold}GPU.cc) + set_source_files_properties(${_cuda_sources} PROPERTIES LANGUAGE HIP) endif() endforeach() @@ -353,6 +356,7 @@ foreach(_bond ${_bonds}) export_PotentialBond${_bond}GPU.cc PotentialBond${_bond}GPUKernel.cu ) + set_source_files_properties(${_cuda_sources} PROPERTIES LANGUAGE HIP) endif() endforeach() @@ -373,6 +377,7 @@ foreach(_bond ${_bonds}) export_PotentialMeshBond${_bond}GPU.cc PotentialMeshBond${_bond}GPUKernel.cu ) + set_source_files_properties(${_cuda_sources} PROPERTIES LANGUAGE HIP) endif() endforeach() @@ -396,6 +401,7 @@ foreach(_pair ${_pairs}) export_PotentialSpecialPair${_pair}GPU.cc PotentialSpecialPair${_pair}GPUKernel.cu ) + set_source_files_properties(${_cuda_sources} PROPERTIES LANGUAGE HIP) endif() endforeach() @@ -418,6 +424,7 @@ foreach(_evaluator ${_triplets}) export_PotentialTersoff${_evaluator}GPU.cc PotentialTersoff${_evaluator}GPUKernel.cu ) + set_source_files_properties(${_cuda_sources} PROPERTIES LANGUAGE HIP) endif() endforeach() @@ -440,6 +447,7 @@ foreach(_evaluator ${_external_evaluators}) export_PotentialExternal${_evaluator}GPU.cc PotentialExternal${_evaluator}GPUKernel.cu ) + set_source_files_properties(${_cuda_sources} PROPERTIES LANGUAGE HIP) endif() endforeach() @@ -462,6 +470,7 @@ foreach(_evaluator ${_wall_evaluators}) export_PotentialExternalWall${_evaluator}GPU.cc PotentialExternalWall${_evaluator}GPUKernel.cu ) + set_source_files_properties(${_cuda_sources} PROPERTIES LANGUAGE HIP) endif() endforeach() @@ -484,6 +493,7 @@ foreach(_evaluator ${_dpdthermo_evaluators}) export_PotentialPairDPDThermo${_evaluator}GPU.cc PotentialPairDPDThermo${_evaluator}GPUKernel.cu ) + set_source_files_properties(${_cuda_sources} PROPERTIES LANGUAGE HIP) endif() endforeach() @@ -532,6 +542,7 @@ foreach(_evaluator ${_pair_evaluators}) export_PotentialPair${_evaluator}GPU.cc PotentialPair${_evaluator}GPUKernel.cu ) + set_source_files_properties(${_cuda_sources} PROPERTIES LANGUAGE HIP) endif() endforeach() From c889bd586913a89ef63295e8ee1aea2f588fcd69 Mon Sep 17 00:00:00 2001 From: Jens Glaser Date: Thu, 16 Jun 2022 09:55:47 -0400 Subject: [PATCH 03/15] fix compile errors w/HIP --- hoomd/md/TwoStepRATTLENVEGPU.cu | 4 ++-- hoomd/md/TwoStepRATTLENVEGPU.cuh | 4 ++-- 2 files changed, 4 insertions(+), 4 deletions(-) diff --git a/hoomd/md/TwoStepRATTLENVEGPU.cu b/hoomd/md/TwoStepRATTLENVEGPU.cu index 50f44636f9..aa626d7f6d 100644 --- a/hoomd/md/TwoStepRATTLENVEGPU.cu +++ b/hoomd/md/TwoStepRATTLENVEGPU.cu @@ -134,8 +134,8 @@ hipError_t gpu_rattle_nve_step_one(Scalar4* d_pos, unsigned int block_size) { unsigned int max_block_size; - cudaFuncAttributes attr; - cudaFuncGetAttributes(&attr, (const void*)gpu_rattle_nve_step_one_kernel); + hipFuncAttributes attr; + hipFuncGetAttributes(&attr, (const void*)gpu_rattle_nve_step_one_kernel); max_block_size = attr.maxThreadsPerBlock; unsigned int run_block_size = min(block_size, max_block_size); diff --git a/hoomd/md/TwoStepRATTLENVEGPU.cuh b/hoomd/md/TwoStepRATTLENVEGPU.cuh index 36d3313e0d..380cebcbd9 100644 --- a/hoomd/md/TwoStepRATTLENVEGPU.cuh +++ b/hoomd/md/TwoStepRATTLENVEGPU.cuh @@ -386,9 +386,9 @@ __global__ void gpu_include_rattle_force_nve_kernel(const Scalar4* d_pos, } while (resid > tolerance && iteration < maxiteration); - accel -= lambda * normal; + accel = accel - lambda * normal; - force -= inv_mass * lambda * normal; + force = force - inv_mass * lambda * normal; virial0 -= lambda * normal.x * pos.x; virial1 -= 0.5 * lambda * (normal.x * pos.y + normal.y * pos.x); From 9af786568ad3ba05b8731ff05bf82af273150781 Mon Sep 17 00:00:00 2001 From: Jens Glaser Date: Mon, 27 Jun 2022 14:49:14 -0400 Subject: [PATCH 04/15] fix more compile errors --- hoomd/md/CMakeLists.txt | 20 ++++++++++---------- 1 file changed, 10 insertions(+), 10 deletions(-) diff --git a/hoomd/md/CMakeLists.txt b/hoomd/md/CMakeLists.txt index 3749b69ff9..e796d53711 100644 --- a/hoomd/md/CMakeLists.txt +++ b/hoomd/md/CMakeLists.txt @@ -311,8 +311,8 @@ foreach(_manifold ${_manifolds}) configure_file(TwoStepRATTLEGPU.cu.inc TwoStepRATTLE${_manifold}GPU.cu @ONLY) + set(_md_sources ${_md_sources} export_ActiveForceConstraintCompute${_manifold}GPU.cc) set(_cuda_sources ${_cuda_sources} - export_ActiveForceConstraintCompute${_manifold}GPU.cc ActiveForceConstraintCompute${_manifold}GPU.cu TwoStepRATTLE${_manifold}GPU.cu ) @@ -329,7 +329,7 @@ foreach(_manifold ${_manifolds}) configure_file(export_TwoStepRATTLEGPU.cc.inc export_TwoStepRATTLE${_method}${_manifold}GPU.cc @ONLY) - set(_cuda_sources ${_cuda_sources} export_TwoStepRATTLE${_method}${_manifold}GPU.cc) + set(_md_sources ${_md_sources} export_TwoStepRATTLE${_method}${_manifold}GPU.cc) set_source_files_properties(${_cuda_sources} PROPERTIES LANGUAGE HIP) endif() endforeach() @@ -352,8 +352,8 @@ foreach(_bond ${_bonds}) configure_file(PotentialBondGPUKernel.cu.inc PotentialBond${_bond}GPUKernel.cu @ONLY) + set(_md_sources ${_md_sources} export_PotentialBond${_bond}GPU.cc) set(_cuda_sources ${_cuda_sources} - export_PotentialBond${_bond}GPU.cc PotentialBond${_bond}GPUKernel.cu ) set_source_files_properties(${_cuda_sources} PROPERTIES LANGUAGE HIP) @@ -373,8 +373,8 @@ foreach(_bond ${_bonds}) configure_file(PotentialMeshBondGPUKernel.cu.inc PotentialMeshBond${_bond}GPUKernel.cu @ONLY) + set(_md_sources ${_md_sources} export_PotentialMeshBond${_bond}GPU.cc) set(_cuda_sources ${_cuda_sources} - export_PotentialMeshBond${_bond}GPU.cc PotentialMeshBond${_bond}GPUKernel.cu ) set_source_files_properties(${_cuda_sources} PROPERTIES LANGUAGE HIP) @@ -397,8 +397,8 @@ foreach(_pair ${_pairs}) configure_file(PotentialSpecialPairGPUKernel.cu.inc PotentialSpecialPair${_pair}GPUKernel.cu @ONLY) + set(_md_sources ${_md_sources} export_PotentialSpecialPair${_pair}GPU.cc) set(_cuda_sources ${_cuda_sources} - export_PotentialSpecialPair${_pair}GPU.cc PotentialSpecialPair${_pair}GPUKernel.cu ) set_source_files_properties(${_cuda_sources} PROPERTIES LANGUAGE HIP) @@ -420,8 +420,8 @@ foreach(_evaluator ${_triplets}) configure_file(PotentialTersoffGPUKernel.cu.inc PotentialTersoff${_evaluator}GPUKernel.cu @ONLY) + set(_md_sources ${_md_sources} export_PotentialTersoff${_evaluator}GPU.cc) set(_cuda_sources ${_cuda_sources} - export_PotentialTersoff${_evaluator}GPU.cc PotentialTersoff${_evaluator}GPUKernel.cu ) set_source_files_properties(${_cuda_sources} PROPERTIES LANGUAGE HIP) @@ -443,8 +443,8 @@ foreach(_evaluator ${_external_evaluators}) configure_file(PotentialExternalGPUKernel.cu.inc PotentialExternal${_evaluator}GPUKernel.cu @ONLY) + set(_md_sources ${_md_sources} export_PotentialExternal${_evaluator}GPU.cc) set(_cuda_sources ${_cuda_sources} - export_PotentialExternal${_evaluator}GPU.cc PotentialExternal${_evaluator}GPUKernel.cu ) set_source_files_properties(${_cuda_sources} PROPERTIES LANGUAGE HIP) @@ -466,8 +466,8 @@ foreach(_evaluator ${_wall_evaluators}) configure_file(PotentialExternalWallGPUKernel.cu.inc PotentialExternalWall${_evaluator}GPUKernel.cu @ONLY) + set(_md_sources ${_md_sources} export_PotentialExternalWall${_evaluator}GPU.cc) set(_cuda_sources ${_cuda_sources} - export_PotentialExternalWall${_evaluator}GPU.cc PotentialExternalWall${_evaluator}GPUKernel.cu ) set_source_files_properties(${_cuda_sources} PROPERTIES LANGUAGE HIP) @@ -489,8 +489,8 @@ foreach(_evaluator ${_dpdthermo_evaluators}) configure_file(PotentialPairDPDThermoGPUKernel.cu.inc PotentialPairDPDThermo${_evaluator}GPUKernel.cu @ONLY) + set(_md_sources ${_md_sources} export_PotentialPairDPDThermo${_evaluator}GPU.cc) set(_cuda_sources ${_cuda_sources} - export_PotentialPairDPDThermo${_evaluator}GPU.cc PotentialPairDPDThermo${_evaluator}GPUKernel.cu ) set_source_files_properties(${_cuda_sources} PROPERTIES LANGUAGE HIP) @@ -538,8 +538,8 @@ foreach(_evaluator ${_pair_evaluators}) configure_file(PotentialPairGPUKernel.cu.inc PotentialPair${_evaluator}GPUKernel.cu @ONLY) + set(_md_sources ${_md_sources} export_PotentialPair${_evaluator}GPU.cc) set(_cuda_sources ${_cuda_sources} - export_PotentialPair${_evaluator}GPU.cc PotentialPair${_evaluator}GPUKernel.cu ) set_source_files_properties(${_cuda_sources} PROPERTIES LANGUAGE HIP) From a3208dbedcf06a78d37b2eaa5f69b71914e12f43 Mon Sep 17 00:00:00 2001 From: Jens Glaser Date: Mon, 27 Jun 2022 14:49:30 -0400 Subject: [PATCH 05/15] use static_cast instead of Cstye conversion --- hoomd/md/EvaluatorTersoff.h | 8 ++++---- hoomd/md/ForceComposite.h | 14 +++++++------- 2 files changed, 11 insertions(+), 11 deletions(-) diff --git a/hoomd/md/EvaluatorTersoff.h b/hoomd/md/EvaluatorTersoff.h index 8611b3befd..2f213a0254 100644 --- a/hoomd/md/EvaluatorTersoff.h +++ b/hoomd/md/EvaluatorTersoff.h @@ -91,15 +91,15 @@ class EvaluatorTersoff pybind11::dict v; pybind11::list mags; - mags.append((Scalar)coeffs.x); - mags.append((Scalar)coeffs.y); + mags.append(static_cast(coeffs.x)); + mags.append(static_cast(coeffs.y)); v["magnitudes"] = pybind11::tuple(mags); v["cutoff_thickness"] = cutoff_thickness; pybind11::list exp_factors; - exp_factors.append((Scalar)exp_consts.x); - exp_factors.append((Scalar)exp_consts.y); + exp_factors.append(static_cast(exp_consts.x)); + exp_factors.append(static_cast(exp_consts.y)); v["exp_factors"] = pybind11::tuple(exp_factors); v["lambda3"] = pow(lambda_cube, 1. / 3.); diff --git a/hoomd/md/ForceComposite.h b/hoomd/md/ForceComposite.h index 7d633386ae..e15beaaa32 100644 --- a/hoomd/md/ForceComposite.h +++ b/hoomd/md/ForceComposite.h @@ -172,13 +172,13 @@ class PYBIND11_EXPORT ForceComposite : public MolecularForceCompute for (unsigned int i = 0; i < N; i++) { auto index = m_body_idx(body_type_id, i); - positions.append(pybind11::make_tuple((Scalar)h_body_pos.data[index].x, - (Scalar)h_body_pos.data[index].y, - (Scalar)h_body_pos.data[index].z)); - orientations.append(pybind11::make_tuple((Scalar)h_body_orientation.data[index].x, - (Scalar)h_body_orientation.data[index].y, - (Scalar)h_body_orientation.data[index].z, - (Scalar)h_body_orientation.data[index].w)); + positions.append(pybind11::make_tuple(static_cast(h_body_pos.data[index].x), + static_cast(h_body_pos.data[index].y), + static_cast(h_body_pos.data[index].z))); + orientations.append(pybind11::make_tuple(static_cast(h_body_orientation.data[index].x), + static_cast(h_body_orientation.data[index].y), + static_cast(h_body_orientation.data[index].z), + static_cast(h_body_orientation.data[index].w))); types.append(m_pdata->getNameByType(h_body_types.data[index])); charges.append(m_body_charge[body_type_id][i]); diameters.append(m_body_diameter[body_type_id][i]); From 554bc0d5cfe30ed9fb0ad4e76487bc42c95c1249 Mon Sep 17 00:00:00 2001 From: Jens Glaser Date: Mon, 27 Jun 2022 14:51:54 -0400 Subject: [PATCH 06/15] remove currently obsolete comments --- CMake/hoomd/FindCUDALibs.cmake | 56 ---------------------------------- 1 file changed, 56 deletions(-) diff --git a/CMake/hoomd/FindCUDALibs.cmake b/CMake/hoomd/FindCUDALibs.cmake index deb96ab125..61bb54973b 100644 --- a/CMake/hoomd/FindCUDALibs.cmake +++ b/CMake/hoomd/FindCUDALibs.cmake @@ -161,64 +161,8 @@ else() add_library(CUDA::cusparse UNKNOWN IMPORTED) endif() - -#find_library(HIP_hipsparse_LIBRARY hipsparse -# PATHS -# "${HIP_ROOT_DIR}" -# ENV ROCM_PATH -# ENV HIP_PATH -# /opt/rocm -# /opt/rocm/hipsparse -# PATH_SUFFIXES lib -# NO_DEFAULT_PATH) -#find_path(HIP_hipsparse_INCLUDE_DIR -# NAMES hipsparse.h -# PATHS -# ${HIP_ROOT_DIR}/hipsparse/include -# $ENV{ROCM_PATH}/hipsparse/include -# $ENV{HIP_PATH}/hipsparse/include -# /opt/rocm/include -# /opt/rocm/hipsparse/include -# NO_DEFAULT_PATH) -#mark_as_advanced(HIP_hipsparse_LIBRARY) -#list(APPEND REQUIRED_CUDA_LIB_VARS HIP_hipsparse_LIBRARY) -#list(APPEND _hipsparse_includes ${HIP_hipsparse_INCLUDE_DIR}) - -#if(HIP_hipsparse_LIBRARY AND NOT TARGET HIP::hipsparse) -# add_library(HIP::hipsparse UNKNOWN IMPORTED) -# set_target_properties(HIP::hipsparse PROPERTIES -# IMPORTED_LOCATION "${HIP_hipsparse_LIBRARY}" -# ) -#endif() - if (HIP_PLATFORM STREQUAL "hip-clang") find_package(hipfft) -# find_library(HIP_rocsparse_LIBRARY rocsparse -# PATHS -# "${HIP_ROOT_DIR}" -# ENV ROCM_PATH -# ENV HIP_PATH -# /opt/rocm -# /opt/rocm/rocsparse -# PATH_SUFFIXES lib -# NO_DEFAULT_PATH) -# find_path(HIP_rocsparse_INCLUDE_DIR -# NAMES rocsparse.h -# PATHS -# ${HIP_ROOT_DIR}/rocsparse/include -# $ENV{ROCM_PATH}/rocsparse/include -# $ENV{HIP_PATH}/rocsparse/include -# /opt/rocm/include -# /opt/rocm/rocsparse/include -# NO_DEFAULT_PATH) -# -# list(APPEND _hipsparse_includes ${HIP_rocsparse_INCLUDE_DIR}) -# mark_as_advanced(HIP_rocsparse_LIBRARY) -# set_target_properties(HIP::hipsparse PROPERTIES -# INTERFACE_INCLUDE_DIRECTORIES "${_hipsparse_includes}" -# INTERFACE_LINK_LIBRARIES "${HIP_rocsparse_LIBRARY}" -# ) -# list(APPEND REQUIRED_CUDA_LIB_VARS HIP_rocsparse_LIBRARY) endif() if (HIP_PLATFORM STREQUAL "nvcc") From 29889074e099b0b9b2c909825ed89b121516c449 Mon Sep 17 00:00:00 2001 From: Jens Glaser Date: Mon, 27 Jun 2022 14:56:21 -0400 Subject: [PATCH 07/15] remove more unneeded code --- CMake/hoomd/HOOMDHIPSetup.cmake | 32 -------------------------------- 1 file changed, 32 deletions(-) diff --git a/CMake/hoomd/HOOMDHIPSetup.cmake b/CMake/hoomd/HOOMDHIPSetup.cmake index 9c14f9a64f..686e0e8ae0 100644 --- a/CMake/hoomd/HOOMDHIPSetup.cmake +++ b/CMake/hoomd/HOOMDHIPSetup.cmake @@ -7,38 +7,6 @@ if(ENABLE_HIP) # setup nvcc to build for all CUDA architectures. Allow user to modify the list if desired set(CMAKE_HIP_ARCHITECTURES gfx900 gfx906 gfx908 gfx90a CACHE STRING "List of AMD GPU to compile HIP code for. Separate with semicolons.") set(HIP_PLATFORM hip-clang) - - #foreach(_amdgpu_target ${AMDGPU_TARGET_LIST}) - # set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} --amdgpu-target=${_amdgpu_target}") - #endforeach (_amdgpu_target) - - #if (HIP_FOUND) - # # reduce link time (no device linking) - # set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} -fno-gpu-rdc") - #endif() - - #search for HIP include directory -# find_path(HIP_INCLUDE_DIR hip/hip_runtime.h -# PATHS -# "${HIP_ROOT_DIR}" -# ENV ROCM_PATH -# ENV HIP_PATH -# PATH_SUFFIXES include) -# -# find_path(ROCm_hsa_INCLUDE_DIR -# NAMES hsa/hsa.h -# PATHS -# ${HIP_ROOT_DIR}/hsa -# ${HIP_ROOT_DIR}/hsa -# $ENV{ROCM_PATH}/hsa -# $ENV{HIP_PATH}/hsa -# $ENV{HSA_PATH} -# /opt/rocm -# PATH_SUFFIXES include -# NO_DEFAULT_PATH) -# -# -# list(APPEND HIP_INCLUDE_DIR ${ROCm_hsa_INCLUDE_DIR}) else() # here we go if hipcc is not available, fall back on internal HIP->CUDA headers ENABLE_LANGUAGE(CUDA) From c7334e676f9a95609f40bc468d72e2b09c931bc2 Mon Sep 17 00:00:00 2001 From: Jens Glaser Date: Mon, 27 Jun 2022 15:13:15 -0400 Subject: [PATCH 08/15] re-enable compilation with nvcc --- CMake/hoomd/HOOMDHIPSetup.cmake | 30 ++++++++++++++++-------------- hoomd/CMakeLists.txt | 2 +- hoomd/md/CMakeLists.txt | 22 +++++++++++----------- 3 files changed, 28 insertions(+), 26 deletions(-) diff --git a/CMake/hoomd/HOOMDHIPSetup.cmake b/CMake/hoomd/HOOMDHIPSetup.cmake index 686e0e8ae0..fe5bab4d20 100644 --- a/CMake/hoomd/HOOMDHIPSetup.cmake +++ b/CMake/hoomd/HOOMDHIPSetup.cmake @@ -3,6 +3,7 @@ if(ENABLE_HIP) if (HIP_FOUND) ENABLE_LANGUAGE(HIP) + SET(HOOMD_DEVICE_LANGUAGE HIP) # setup nvcc to build for all CUDA architectures. Allow user to modify the list if desired set(CMAKE_HIP_ARCHITECTURES gfx900 gfx906 gfx908 gfx90a CACHE STRING "List of AMD GPU to compile HIP code for. Separate with semicolons.") @@ -10,6 +11,7 @@ if(ENABLE_HIP) else() # here we go if hipcc is not available, fall back on internal HIP->CUDA headers ENABLE_LANGUAGE(CUDA) + SET(HOOMD_DEVICE_LANGUAGE CUDA) set(HIP_INCLUDE_DIR "$,${CMAKE_CURRENT_SOURCE_DIR},${HOOMD_INSTALL_PREFIX}/${PYTHON_SITE_INSTALL_DIR}/include>/hoomd/extern/HIP/include/") @@ -32,19 +34,19 @@ if(ENABLE_HIP) endif() endif() -# if(NOT TARGET HIP::hip) -# add_library(HIP::hip INTERFACE IMPORTED) -# set_target_properties(HIP::hip PROPERTIES -# INTERFACE_INCLUDE_DIRECTORIES "${HIP_INCLUDE_DIR};${HIPCUB_INCLUDE_DIR}") -# -# # set HIP_VERSION_* on non-CUDA targets (the version is already defined on AMD targets through hipcc) -# set_property(TARGET HIP::hip APPEND PROPERTY INTERFACE_COMPILE_DEFINITIONS -# $<$>:HIP_VERSION_MAJOR=${HIP_VERSION_MAJOR}>) -# set_property(TARGET HIP::hip APPEND PROPERTY INTERFACE_COMPILE_DEFINITIONS -# $<$>:HIP_VERSION_MINOR=${HIP_VERSION_MINOR}>) -# set_property(TARGET HIP::hip APPEND PROPERTY INTERFACE_COMPILE_DEFINITIONS -# $<$>:HIP_VERSION_PATCH=${HIP_VERSION_PATCH}>) -# + if(NOT TARGET hip::host) + add_library(hip::host INTERFACE IMPORTED) + set_target_properties(hip::host PROPERTIES + INTERFACE_INCLUDE_DIRECTORIES "${HIP_INCLUDE_DIR};${HIPCUB_INCLUDE_DIR}") + + # set HIP_VERSION_* on non-CUDA targets (the version is already defined on AMD targets through hipcc) + set_property(TARGET hip::host APPEND PROPERTY INTERFACE_COMPILE_DEFINITIONS + $<$>:HIP_VERSION_MAJOR=${HIP_VERSION_MAJOR}>) + set_property(TARGET hip::host APPEND PROPERTY INTERFACE_COMPILE_DEFINITIONS + $<$>:HIP_VERSION_MINOR=${HIP_VERSION_MINOR}>) + set_property(TARGET hip::host APPEND PROPERTY INTERFACE_COMPILE_DEFINITIONS + $<$>:HIP_VERSION_PATCH=${HIP_VERSION_PATCH}>) + # # branch upon HCC or NVCC target # if(${HIP_PLATFORM} STREQUAL "nvcc") # set_property(TARGET HIP::hip APPEND PROPERTY INTERFACE_COMPILE_DEFINITIONS @@ -53,7 +55,7 @@ if(ENABLE_HIP) # set_property(TARGET HIP::hip APPEND PROPERTY INTERFACE_COMPILE_DEFINITIONS # $<$>:__HIP_PLATFORM_HCC__>) # endif() -# endif() + endif() # branch upon HCC or NVCC target if(${HIP_PLATFORM} STREQUAL "nvcc") diff --git a/hoomd/CMakeLists.txt b/hoomd/CMakeLists.txt index a345309de9..61efc290d7 100644 --- a/hoomd/CMakeLists.txt +++ b/hoomd/CMakeLists.txt @@ -212,7 +212,7 @@ set(_hoomd_cu_sources BondedGroupData.cu if (ENABLE_HIP) set(_cuda_sources ${_hoomd_cu_sources}) -set_source_files_properties(${_hoomd_cu_sources} PROPERTIES LANGUAGE HIP) +set_source_files_properties(${_hoomd_cu_sources} PROPERTIES LANGUAGE ${HOOMD_DEVICE_LANGUAGE}) endif (ENABLE_HIP) ######################### diff --git a/hoomd/md/CMakeLists.txt b/hoomd/md/CMakeLists.txt index e796d53711..327a658cd7 100644 --- a/hoomd/md/CMakeLists.txt +++ b/hoomd/md/CMakeLists.txt @@ -274,7 +274,7 @@ set(_md_cu_sources ActiveForceComputeGPU.cu if (ENABLE_HIP) set(_cuda_sources ${_md_cu_sources} ${DFFT_CU_SOURCES}) -set_source_files_properties(${_md_cu_sources} PROPERTIES LANGUAGE HIP) +set_source_files_properties(${_md_cu_sources} PROPERTIES LANGUAGE ${HOOMD_DEVICE_LANGUAGE}) endif (ENABLE_HIP) # generate pybind11 export cc files @@ -316,7 +316,7 @@ foreach(_manifold ${_manifolds}) ActiveForceConstraintCompute${_manifold}GPU.cu TwoStepRATTLE${_manifold}GPU.cu ) - set_source_files_properties(${_cuda_sources} PROPERTIES LANGUAGE HIP) + set_source_files_properties(${_cuda_sources} PROPERTIES LANGUAGE ${HOOMD_DEVICE_LANGUAGE}) endif() foreach (_method ${_rattle_methods}) @@ -330,7 +330,7 @@ foreach(_manifold ${_manifolds}) export_TwoStepRATTLE${_method}${_manifold}GPU.cc @ONLY) set(_md_sources ${_md_sources} export_TwoStepRATTLE${_method}${_manifold}GPU.cc) - set_source_files_properties(${_cuda_sources} PROPERTIES LANGUAGE HIP) + set_source_files_properties(${_cuda_sources} PROPERTIES LANGUAGE ${HOOMD_DEVICE_LANGUAGE}) endif() endforeach() @@ -356,7 +356,7 @@ foreach(_bond ${_bonds}) set(_cuda_sources ${_cuda_sources} PotentialBond${_bond}GPUKernel.cu ) - set_source_files_properties(${_cuda_sources} PROPERTIES LANGUAGE HIP) + set_source_files_properties(${_cuda_sources} PROPERTIES LANGUAGE ${HOOMD_DEVICE_LANGUAGE}) endif() endforeach() @@ -377,7 +377,7 @@ foreach(_bond ${_bonds}) set(_cuda_sources ${_cuda_sources} PotentialMeshBond${_bond}GPUKernel.cu ) - set_source_files_properties(${_cuda_sources} PROPERTIES LANGUAGE HIP) + set_source_files_properties(${_cuda_sources} PROPERTIES LANGUAGE ${HOOMD_DEVICE_LANGUAGE}) endif() endforeach() @@ -401,7 +401,7 @@ foreach(_pair ${_pairs}) set(_cuda_sources ${_cuda_sources} PotentialSpecialPair${_pair}GPUKernel.cu ) - set_source_files_properties(${_cuda_sources} PROPERTIES LANGUAGE HIP) + set_source_files_properties(${_cuda_sources} PROPERTIES LANGUAGE ${HOOMD_DEVICE_LANGUAGE}) endif() endforeach() @@ -424,7 +424,7 @@ foreach(_evaluator ${_triplets}) set(_cuda_sources ${_cuda_sources} PotentialTersoff${_evaluator}GPUKernel.cu ) - set_source_files_properties(${_cuda_sources} PROPERTIES LANGUAGE HIP) + set_source_files_properties(${_cuda_sources} PROPERTIES LANGUAGE ${HOOMD_DEVICE_LANGUAGE}) endif() endforeach() @@ -447,7 +447,7 @@ foreach(_evaluator ${_external_evaluators}) set(_cuda_sources ${_cuda_sources} PotentialExternal${_evaluator}GPUKernel.cu ) - set_source_files_properties(${_cuda_sources} PROPERTIES LANGUAGE HIP) + set_source_files_properties(${_cuda_sources} PROPERTIES LANGUAGE ${HOOMD_DEVICE_LANGUAGE}) endif() endforeach() @@ -470,7 +470,7 @@ foreach(_evaluator ${_wall_evaluators}) set(_cuda_sources ${_cuda_sources} PotentialExternalWall${_evaluator}GPUKernel.cu ) - set_source_files_properties(${_cuda_sources} PROPERTIES LANGUAGE HIP) + set_source_files_properties(${_cuda_sources} PROPERTIES LANGUAGE ${HOOMD_DEVICE_LANGUAGE}) endif() endforeach() @@ -493,7 +493,7 @@ foreach(_evaluator ${_dpdthermo_evaluators}) set(_cuda_sources ${_cuda_sources} PotentialPairDPDThermo${_evaluator}GPUKernel.cu ) - set_source_files_properties(${_cuda_sources} PROPERTIES LANGUAGE HIP) + set_source_files_properties(${_cuda_sources} PROPERTIES LANGUAGE ${HOOMD_DEVICE_LANGUAGE}) endif() endforeach() @@ -542,7 +542,7 @@ foreach(_evaluator ${_pair_evaluators}) set(_cuda_sources ${_cuda_sources} PotentialPair${_evaluator}GPUKernel.cu ) - set_source_files_properties(${_cuda_sources} PROPERTIES LANGUAGE HIP) + set_source_files_properties(${_cuda_sources} PROPERTIES LANGUAGE ${HOOMD_DEVICE_LANGUAGE}) endif() endforeach() From 7b62cb674c5ee95f6bc3a0514a7167cc0ee90f09 Mon Sep 17 00:00:00 2001 From: jglaser Date: Tue, 19 Jul 2022 08:36:59 -0400 Subject: [PATCH 09/15] Update hoomd/md/OPLSDihedralForceCompute.cc Co-authored-by: Joshua A. Anderson --- hoomd/md/OPLSDihedralForceCompute.cc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/hoomd/md/OPLSDihedralForceCompute.cc b/hoomd/md/OPLSDihedralForceCompute.cc index bb366743e5..fc8e45194b 100644 --- a/hoomd/md/OPLSDihedralForceCompute.cc +++ b/hoomd/md/OPLSDihedralForceCompute.cc @@ -338,7 +338,7 @@ void OPLSDihedralForceCompute::computeForces(uint64_t timestep) h_force.data[i4].x = h_force.data[i4].x + f4.x; h_force.data[i4].y = h_force.data[i4].y + f4.y; h_force.data[i4].z = h_force.data[i4].z + f4.z; - h_force.data[i4].w = h_force.data[i4].z + f4.w; + h_force.data[i4].w = h_force.data[i4].w + f4.w; // Compute 1/4 of the virial, 1/4 for each atom in the dihedral // upper triangular version of virial tensor From 5b2f5803981d0fa0a0f47733edc046f52de7eef3 Mon Sep 17 00:00:00 2001 From: "pre-commit-ci[bot]" <66853113+pre-commit-ci[bot]@users.noreply.github.com> Date: Tue, 19 Jul 2022 16:22:55 +0000 Subject: [PATCH 10/15] [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci --- hoomd/md/ForceComposite.h | 9 +++++---- hoomd/md/TwoStepRATTLEBDGPU.h | 8 ++++---- 2 files changed, 9 insertions(+), 8 deletions(-) diff --git a/hoomd/md/ForceComposite.h b/hoomd/md/ForceComposite.h index e15beaaa32..98c792c74c 100644 --- a/hoomd/md/ForceComposite.h +++ b/hoomd/md/ForceComposite.h @@ -175,10 +175,11 @@ class PYBIND11_EXPORT ForceComposite : public MolecularForceCompute positions.append(pybind11::make_tuple(static_cast(h_body_pos.data[index].x), static_cast(h_body_pos.data[index].y), static_cast(h_body_pos.data[index].z))); - orientations.append(pybind11::make_tuple(static_cast(h_body_orientation.data[index].x), - static_cast(h_body_orientation.data[index].y), - static_cast(h_body_orientation.data[index].z), - static_cast(h_body_orientation.data[index].w))); + orientations.append( + pybind11::make_tuple(static_cast(h_body_orientation.data[index].x), + static_cast(h_body_orientation.data[index].y), + static_cast(h_body_orientation.data[index].z), + static_cast(h_body_orientation.data[index].w))); types.append(m_pdata->getNameByType(h_body_types.data[index])); charges.append(m_body_charge[body_type_id][i]); diameters.append(m_body_diameter[body_type_id][i]); diff --git a/hoomd/md/TwoStepRATTLEBDGPU.h b/hoomd/md/TwoStepRATTLEBDGPU.h index f9bc4a27a1..dc60bb983d 100644 --- a/hoomd/md/TwoStepRATTLEBDGPU.h +++ b/hoomd/md/TwoStepRATTLEBDGPU.h @@ -144,7 +144,7 @@ template void TwoStepRATTLEBDGPU::integrateStepOne(uin bool aniso = this->m_aniso; - #if defined(__HIP_PLATFORM_NVCC__) +#if defined(__HIP_PLATFORM_NVCC__) if (this->m_exec_conf->allConcurrentManagedAccess()) { // prefetch gammas @@ -161,7 +161,7 @@ template void TwoStepRATTLEBDGPU::integrateStepOne(uin if (this->m_exec_conf->isCUDAErrorCheckingEnabled()) CHECK_CUDA_ERROR(); } - #endif +#endif this->m_exec_conf->beginMultiGPU(); @@ -233,7 +233,7 @@ template void TwoStepRATTLEBDGPU::includeRATTLEForce(u args.timestep = timestep; args.seed = this->m_sysdef->getSeed(); - #if defined(__HIP_PLATFORM_NVCC__) +#if defined(__HIP_PLATFORM_NVCC__) if (this->m_exec_conf->allConcurrentManagedAccess()) { // prefetch gammas @@ -247,7 +247,7 @@ template void TwoStepRATTLEBDGPU::includeRATTLEForce(u if (this->m_exec_conf->isCUDAErrorCheckingEnabled()) CHECK_CUDA_ERROR(); } - #endif +#endif this->m_exec_conf->beginMultiGPU(); From ba337738bfdf67878814b2b7acb9965f970258ab Mon Sep 17 00:00:00 2001 From: "Joshua A. Anderson" Date: Thu, 21 Jul 2022 10:12:01 -0400 Subject: [PATCH 11/15] Suppress HIP not found errors on CUDA systems. --- CMake/hoomd/HOOMDHIPSetup.cmake | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/CMake/hoomd/HOOMDHIPSetup.cmake b/CMake/hoomd/HOOMDHIPSetup.cmake index fe5bab4d20..97e1dfac08 100644 --- a/CMake/hoomd/HOOMDHIPSetup.cmake +++ b/CMake/hoomd/HOOMDHIPSetup.cmake @@ -1,7 +1,8 @@ if(ENABLE_HIP) - find_package(HIP) + find_package(HIP QUIET) if (HIP_FOUND) + find_package(HIP) ENABLE_LANGUAGE(HIP) SET(HOOMD_DEVICE_LANGUAGE HIP) From 99d4f53f902b1f86096e022289e75d12deee68cf Mon Sep 17 00:00:00 2001 From: "Joshua A. Anderson" Date: Thu, 21 Jul 2022 10:16:48 -0400 Subject: [PATCH 12/15] Document minimum CMake requirement for HIP. --- BUILDING.rst | 6 ++++++ CMake/hoomd/HOOMDHIPSetup.cmake | 1 + 2 files changed, 7 insertions(+) diff --git a/BUILDING.rst b/BUILDING.rst index 8e6175ba7d..d5df98ade3 100644 --- a/BUILDING.rst +++ b/BUILDING.rst @@ -105,6 +105,7 @@ Install prerequisites external dependency when building for AMD GPUs - roctracer-dev - Linux kernel >= 3.5.0 + - CMake >= 3.21 For **HOOMD-blue** on AMD GPUs, the following limitations currently apply. @@ -112,6 +113,11 @@ Install prerequisites 2. The ``mpcd`` component is disabled on AMD GPUs. 3. Multi-GPU execution via unified memory is not available. +.. note:: + + When ``ENABLE_GPU=on``, HOOMD-blue will attempt to find HIP first. If it cannot find HIP, it + will attemp to find CUDA. + **For threaded parallelism on the CPU** (required when ``ENABLE_TBB=on``): - Intel Threading Building Blocks >= 4.3 diff --git a/CMake/hoomd/HOOMDHIPSetup.cmake b/CMake/hoomd/HOOMDHIPSetup.cmake index 97e1dfac08..891f4ed03e 100644 --- a/CMake/hoomd/HOOMDHIPSetup.cmake +++ b/CMake/hoomd/HOOMDHIPSetup.cmake @@ -3,6 +3,7 @@ if(ENABLE_HIP) if (HIP_FOUND) find_package(HIP) + CMAKE_MINIMUM_REQUIRED(VERSION 3.21 FATAL_ERROR) ENABLE_LANGUAGE(HIP) SET(HOOMD_DEVICE_LANGUAGE HIP) From 6d8c6a41712642ac8169e13315cc37eb6f84e114 Mon Sep 17 00:00:00 2001 From: "Joshua A. Anderson" Date: Thu, 21 Jul 2022 10:42:38 -0400 Subject: [PATCH 13/15] Add HOOMD_GPU_PLATFORM CMake variable. Allow the user to choose between CUDA and HIP. --- BUILDING.rst | 5 +++-- CMake/hoomd/HOOMDHIPSetup.cmake | 9 +++++---- CMakeLists.txt | 1 + 3 files changed, 9 insertions(+), 6 deletions(-) diff --git a/BUILDING.rst b/BUILDING.rst index d5df98ade3..676af24bf0 100644 --- a/BUILDING.rst +++ b/BUILDING.rst @@ -115,8 +115,8 @@ Install prerequisites .. note:: - When ``ENABLE_GPU=on``, HOOMD-blue will attempt to find HIP first. If it cannot find HIP, it - will attemp to find CUDA. + When ``ENABLE_GPU=on``, HOOMD-blue will default to CUDA. Set ``HHOOMD_GPU_PLATFORM=HIP`` to + choose HIP. **For threaded parallelism on the CPU** (required when ``ENABLE_TBB=on``): @@ -233,6 +233,7 @@ Other option changes take effect at any time: - ``CMAKE_INSTALL_PREFIX`` - Directory to install **HOOMD-blue**. Defaults to the root path of the found Python executable. - ``ENABLE_GPU`` - When enabled, compiled GPU accelerated computations (default: ``off``). +- ``HOOMD_GPU_PLATFORM`` - Choose either ``CUDA`` or ``HIP`` as a GPU backend (default: ``CUDA``). - ``SINGLE_PRECISION`` - Controls precision (default: ``off``). - When set to ``on``, all calculations are performed in single precision. diff --git a/CMake/hoomd/HOOMDHIPSetup.cmake b/CMake/hoomd/HOOMDHIPSetup.cmake index 891f4ed03e..559e4fea94 100644 --- a/CMake/hoomd/HOOMDHIPSetup.cmake +++ b/CMake/hoomd/HOOMDHIPSetup.cmake @@ -1,8 +1,7 @@ if(ENABLE_HIP) - find_package(HIP QUIET) - if (HIP_FOUND) - find_package(HIP) + if (HOOMD_GPU_PLATFORM STREQUAL "HIP") + find_package(HIP REQUIRED) CMAKE_MINIMUM_REQUIRED(VERSION 3.21 FATAL_ERROR) ENABLE_LANGUAGE(HIP) SET(HOOMD_DEVICE_LANGUAGE HIP) @@ -10,7 +9,7 @@ if(ENABLE_HIP) # setup nvcc to build for all CUDA architectures. Allow user to modify the list if desired set(CMAKE_HIP_ARCHITECTURES gfx900 gfx906 gfx908 gfx90a CACHE STRING "List of AMD GPU to compile HIP code for. Separate with semicolons.") set(HIP_PLATFORM hip-clang) - else() + elseif (HOOMD_GPU_PLATFORM STREQUAL "CUDA") # here we go if hipcc is not available, fall back on internal HIP->CUDA headers ENABLE_LANGUAGE(CUDA) SET(HOOMD_DEVICE_LANGUAGE CUDA) @@ -34,6 +33,8 @@ if(ENABLE_HIP) # Use system provided CUB for CUDA 11 and newer set(HIPCUB_INCLUDE_DIR "$,${CMAKE_CURRENT_SOURCE_DIR},${HOOMD_INSTALL_PREFIX}/${PYTHON_SITE_INSTALL_DIR}/include>/hoomd/extern/hipCUB/hipcub/include/") endif() + else() + message(FATAL_ERROR "HOOMD_GPU_PLATFORM must be either CUDA or HIP") endif() if(NOT TARGET hip::host) diff --git a/CMakeLists.txt b/CMakeLists.txt index 812a174b61..3cc866c7f9 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -86,6 +86,7 @@ endif() option(SINGLE_PRECISION "Use single precision math" OFF) OPTION(ENABLE_GPU "True if we are compiling for a GPU target" FALSE) SET(ENABLE_HIP ${ENABLE_GPU}) +set(HOOMD_GPU_PLATFORM "CUDA" CACHE STRING "Choose the GPU backend: HIP or CUDA.") option(ENABLE_LLVM "Link to the LLVM library for run time code generation" off) option(ENABLE_HPMC_MIXED_PRECISION "Enable mixed precision computations in HPMC" ON) From 9011491701295e83a93c525b3c7e7207ad86201d Mon Sep 17 00:00:00 2001 From: jglaser Date: Fri, 29 Jul 2022 11:47:23 -0400 Subject: [PATCH 14/15] Update hoomd/md/BondTablePotential.cc Co-authored-by: Brandon Butler --- hoomd/md/BondTablePotential.cc | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/hoomd/md/BondTablePotential.cc b/hoomd/md/BondTablePotential.cc index 6e99279143..8efa9771ba 100644 --- a/hoomd/md/BondTablePotential.cc +++ b/hoomd/md/BondTablePotential.cc @@ -128,8 +128,8 @@ pybind11::dict BondTablePotential::getParams(std::string type) auto type_id = m_bond_data->getTypeByName(type); pybind11::dict params; - params["r_min"] = (Scalar)h_params.data[type_id].x; - params["r_max"] = (Scalar)h_params.data[type_id].y; + params["r_min"] = static_cast(h_params.data[type_id].x); + params["r_max"] = static_cast(h_params.data[type_id].y); auto V = pybind11::array_t(m_table_width); auto V_unchecked = V.mutable_unchecked<1>(); From b7697ac4f768cf524d04f33b12084b1ec111230d Mon Sep 17 00:00:00 2001 From: jglaser Date: Fri, 29 Jul 2022 11:47:31 -0400 Subject: [PATCH 15/15] Update CMake/hoomd/HOOMDHIPSetup.cmake Co-authored-by: Brandon Butler --- CMake/hoomd/HOOMDHIPSetup.cmake | 8 -------- 1 file changed, 8 deletions(-) diff --git a/CMake/hoomd/HOOMDHIPSetup.cmake b/CMake/hoomd/HOOMDHIPSetup.cmake index 559e4fea94..20305a70a3 100644 --- a/CMake/hoomd/HOOMDHIPSetup.cmake +++ b/CMake/hoomd/HOOMDHIPSetup.cmake @@ -50,14 +50,6 @@ if(ENABLE_HIP) set_property(TARGET hip::host APPEND PROPERTY INTERFACE_COMPILE_DEFINITIONS $<$>:HIP_VERSION_PATCH=${HIP_VERSION_PATCH}>) -# # branch upon HCC or NVCC target -# if(${HIP_PLATFORM} STREQUAL "nvcc") -# set_property(TARGET HIP::hip APPEND PROPERTY INTERFACE_COMPILE_DEFINITIONS -# $<$>:__HIP_PLATFORM_NVCC__>) -# elseif(${HIP_PLATFORM} STREQUAL "hip-clang") -# set_property(TARGET HIP::hip APPEND PROPERTY INTERFACE_COMPILE_DEFINITIONS -# $<$>:__HIP_PLATFORM_HCC__>) -# endif() endif() # branch upon HCC or NVCC target