SWDEV-290637 - Migrate cmake build from hip to hipamd
Change-Id: I605fc633ba18df4fca95ceee83670600dd3f0b71
This commit is contained in:
Anusha GodavarthySurya
committed by
Anusha Godavarthy Surya
Anusha Godavarthy Surya
parent
19a30de0a6
commit
e420815e75
Executable
+345
@@ -0,0 +1,345 @@
|
||||
# Copyright (c) 2020-2021 Advanced Micro Devices, Inc. All rights reserved.
|
||||
#
|
||||
# Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
# of this software and associated documentation files (the "Software"), to deal
|
||||
# in the Software without restriction, including without limitation the rights
|
||||
# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
# copies of the Software, and to permit persons to whom the Software is
|
||||
# furnished to do so, subject to the following conditions:
|
||||
#
|
||||
# The above copyright notice and this permission notice shall be included in
|
||||
# all copies or substantial portions of the Software.
|
||||
#
|
||||
# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
# THE SOFTWARE.
|
||||
|
||||
cmake_minimum_required(VERSION 3.5.1)
|
||||
|
||||
include(GNUInstallDirs)
|
||||
|
||||
if(ADDRESS_SANITIZER)
|
||||
set(ASAN_LINKER_FLAGS "-fsanitize=address")
|
||||
set(ASAN_COMPILER_FLAGS "-fno-omit-frame-pointer -fsanitize=address")
|
||||
|
||||
if(NOT CMAKE_COMPILER_IS_GNUCC)
|
||||
if(BUILD_SHARED_LIBS)
|
||||
set(ASAN_LINKER_FLAGS "${ASAN_LINKER_FLAGS} -shared-libsan")
|
||||
else()
|
||||
set(ASAN_LINKER_FLAGS "${ASAN_LINKER_FLAGS} -static-libsan")
|
||||
endif()
|
||||
endif()
|
||||
|
||||
set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS} ${ASAN_COMPILER_FLAGS}")
|
||||
set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${ASAN_COMPILER_FLAGS}")
|
||||
set(CMAKE_EXE_LINKER_FLAGS "${CMAKE_EXE_LINKER_FLAGS} ${ASAN_LINKER_FLAGS} -s")
|
||||
set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} ${ASAN_LINKER_FLAGS}")
|
||||
endif()
|
||||
|
||||
option(BUILD_SHARED_LIBS "Build the shared library" ON)
|
||||
|
||||
list(APPEND CMAKE_MODULE_PATH "${CMAKE_CURRENT_LIST_DIR}/cmake")
|
||||
find_package(ROCclr)
|
||||
|
||||
if(BUILD_SHARED_LIBS)
|
||||
add_library(amdhip64 SHARED)
|
||||
# Windows doesn't have a strip utility, so CMAKE_STRIP won't be set.
|
||||
if((CMAKE_BUILD_TYPE STREQUAL "Release") AND NOT ("${CMAKE_STRIP}" STREQUAL ""))
|
||||
add_custom_command(TARGET amdhip64 POST_BUILD COMMAND ${CMAKE_STRIP} $<TARGET_FILE:amdhip64>)
|
||||
endif()
|
||||
else()
|
||||
add_library(amdhip64 STATIC $<TARGET_OBJECTS:rocclr>)
|
||||
endif()
|
||||
|
||||
set_target_properties(amdhip64 PROPERTIES
|
||||
CXX_STANDARD 14
|
||||
CXX_STANDARD_REQUIRED ON
|
||||
CXX_EXTENSIONS OFF
|
||||
POSITION_INDEPENDENT_CODE ON
|
||||
# Workaround for many places in the HIP project
|
||||
# having hardcoded references to build/lib/libamdhip64.so
|
||||
LIBRARY_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR}/lib)
|
||||
|
||||
if(CMAKE_SIZEOF_VOID_P EQUAL 8)
|
||||
set_target_properties(amdhip64 PROPERTIES OUTPUT_NAME "amdhip64")
|
||||
else()
|
||||
set_target_properties(amdhip64 PROPERTIES OUTPUT_NAME "amdhip32")
|
||||
endif()
|
||||
|
||||
# Disable versioning for Windows
|
||||
# as currently HIP_LIB_VERSION_STRING and HIP_LIB_VERSION_MAJOR
|
||||
# are not being populated
|
||||
if(NOT WIN32)
|
||||
if(BUILD_SHARED_LIBS)
|
||||
set_target_properties(amdhip64 PROPERTIES
|
||||
VERSION ${HIP_LIB_VERSION_STRING}
|
||||
SOVERSION ${HIP_LIB_VERSION_MAJOR})
|
||||
endif()
|
||||
endif()
|
||||
|
||||
target_sources(amdhip64 PRIVATE
|
||||
cl_gl.cpp
|
||||
cl_lqdflash_amd.cpp
|
||||
fixme.cpp
|
||||
hip_activity.cpp
|
||||
hip_code_object.cpp
|
||||
hip_context.cpp
|
||||
hip_device_runtime.cpp
|
||||
hip_device.cpp
|
||||
hip_error.cpp
|
||||
hip_event.cpp
|
||||
hip_fatbin.cpp
|
||||
hip_global.cpp
|
||||
hip_graph_internal.cpp
|
||||
hip_graph.cpp
|
||||
hip_hmm.cpp
|
||||
hip_intercept.cpp
|
||||
hip_memory.cpp
|
||||
hip_module.cpp
|
||||
hip_peer.cpp
|
||||
hip_platform.cpp
|
||||
hip_profile.cpp
|
||||
hip_rtc.cpp
|
||||
hip_stream_ops.cpp
|
||||
hip_stream.cpp
|
||||
hip_surface.cpp
|
||||
hip_texture.cpp)
|
||||
|
||||
if(WIN32)
|
||||
target_sources(amdhip64 PRIVATE
|
||||
cl_d3d9.cpp
|
||||
cl_d3d10.cpp
|
||||
cl_d3d11.cpp)
|
||||
endif()
|
||||
|
||||
if(BUILD_SHARED_LIBS)
|
||||
if(WIN32)
|
||||
target_sources(amdhip64 PRIVATE amdhip.def)
|
||||
else()
|
||||
target_link_libraries(amdhip64 PRIVATE "-Wl,--version-script=${CMAKE_CURRENT_LIST_DIR}/hip_hcc.map.in")
|
||||
set_target_properties(amdhip64 PROPERTIES LINK_DEPENDS "${CMAKE_CURRENT_LIST_DIR}/hip_hcc.map.in")
|
||||
endif()
|
||||
endif()
|
||||
|
||||
target_include_directories(amdhip64
|
||||
PRIVATE
|
||||
${HIP_COMMON_INCLUDE_DIR}
|
||||
${PROJECT_SOURCE_DIR}/include
|
||||
${PROJECT_BINARY_DIR}/include)
|
||||
|
||||
target_compile_definitions(amdhip64 PRIVATE __HIP_PLATFORM_AMD__)
|
||||
|
||||
target_link_libraries(amdhip64 PRIVATE ${CMAKE_DL_LIBS})
|
||||
# Additional dependencies for hipRTC
|
||||
if(WIN32)
|
||||
target_link_libraries(amdhip64 PRIVATE Dbghelp.lib)
|
||||
endif()
|
||||
|
||||
# Note in static case we cannot link against rocclr.
|
||||
# If we would, we'd also have to export rocclr and have hipcc pass it to the linker.
|
||||
if(BUILD_SHARED_LIBS)
|
||||
target_link_libraries(amdhip64 PRIVATE rocclr)
|
||||
else()
|
||||
target_compile_definitions(amdhip64 PRIVATE $<TARGET_PROPERTY:rocclr,COMPILE_DEFINITIONS>)
|
||||
target_include_directories(amdhip64 PRIVATE $<TARGET_PROPERTY:rocclr,INCLUDE_DIRECTORIES>)
|
||||
endif()
|
||||
|
||||
# Short-Term solution for pre-compiled headers for online compilation
|
||||
# Enable pre compiled header
|
||||
if(__HIP_ENABLE_PCH)
|
||||
find_package(LLVM REQUIRED CONFIG
|
||||
PATHS
|
||||
/opt/rocm/llvm)
|
||||
# find_package(LLVM) returns the lib/cmake/llvm location. We require the root.
|
||||
if(NOT DEFINED HIP_LLVM_ROOT)
|
||||
set(HIP_LLVM_ROOT "${LLVM_DIR}/../../..")
|
||||
endif()
|
||||
|
||||
execute_process(COMMAND sh -c "${HIP_COMMON_BIN_DIR}/hip_embed_pch.sh ${HIP_COMMON_INCLUDE_DIR} ${PROJECT_BINARY_DIR}/include ${PROJECT_SOURCE_DIR}/include ${HIP_LLVM_ROOT}" COMMAND_ECHO STDERR RESULT_VARIABLE EMBED_PCH_RC)
|
||||
if (EMBED_PCH_RC AND NOT EMBED_PCH_RC EQUAL 0)
|
||||
message(FATAL_ERROR "Failed to embed PCH")
|
||||
endif()
|
||||
|
||||
target_compile_definitions(amdhip64 PRIVATE __HIP_ENABLE_PCH)
|
||||
target_sources(amdhip64 PRIVATE ${CMAKE_BINARY_DIR}/hip_pch.o)
|
||||
endif()
|
||||
|
||||
# Enable preprocessed hiprtc-builtins library
|
||||
if(__HIP_ENABLE_RTC)
|
||||
message(STATUS "HIP RTC enabled.")
|
||||
include(HIPRTC RESULT_VARIABLE HIPRTC_CMAKE)
|
||||
# Requires clang and llvm-mc to create this library.
|
||||
find_package(LLVM REQUIRED CONFIG PATHS /opt/rocm/llvm)
|
||||
find_package(Clang REQUIRED CONFIG PATHS /opt/rocm/llvm)
|
||||
set(HIPRTC_GEN_DIR "${CMAKE_CURRENT_BINARY_DIR}/hip_rtc_gen")
|
||||
set(HIPRTC_GEN_HEADER "${HIPRTC_GEN_DIR}/hipRTC_header.h")
|
||||
set(HIPRTC_GEN_MCIN "${HIPRTC_GEN_DIR}/hipRTC_header.mcin")
|
||||
set(HIPRTC_GEN_PREPROCESSED "${HIPRTC_GEN_DIR}/hipRTC")
|
||||
set(HIPRTC_GEN_OBJ "${HIPRTC_GEN_DIR}/hipRTC_header${CMAKE_CXX_OUTPUT_EXTENSION}")
|
||||
|
||||
# Generate required HIPRTC files.
|
||||
FILE(MAKE_DIRECTORY ${HIPRTC_GEN_DIR})
|
||||
generate_hiprtc_header("${HIPRTC_GEN_HEADER}")
|
||||
generate_hiprtc_mcin("${HIPRTC_GEN_MCIN}" "${HIPRTC_GEN_PREPROCESSED}")
|
||||
|
||||
# Generate HIPRTC Builtins Preprocessed Object.
|
||||
# Note: second command appends define macros at build time.
|
||||
# FIXME: --hip-version forced to 3.6 to use clang headers, until Windows versioning is fixed.
|
||||
add_custom_command(
|
||||
OUTPUT ${HIPRTC_GEN_PREPROCESSED}
|
||||
COMMAND $<TARGET_FILE:clang> -O3 --rocm-path=${PROJECT_SOURCE_DIR}/include/.. -std=c++14 -nogpulib --hip-version=3.6 -isystem ${HIP_COMMON_INCLUDE_DIR} -isystem ${PROJECT_SOURCE_DIR}/include -isystem ${PROJECT_BINARY_DIR}/include -isystem ${CMAKE_CURRENT_SOURCE_DIR}/include --cuda-device-only -D__HIPCC_RTC__ -x hip ${HIPRTC_GEN_HEADER} -E -o ${HIPRTC_GEN_PREPROCESSED}
|
||||
COMMAND ${CMAKE_COMMAND} -DHIPRTC_ADD_MACROS=1 -DHIPRTC_PREPROCESSED_FILE=${HIPRTC_GEN_PREPROCESSED} -P ${HIPRTC_CMAKE}
|
||||
DEPENDS clang ${HIPRTC_GEN_HEADER})
|
||||
add_custom_command(
|
||||
OUTPUT ${HIPRTC_GEN_OBJ}
|
||||
COMMAND $<TARGET_FILE:llvm-mc> -o ${HIPRTC_GEN_OBJ} ${HIPRTC_GEN_MCIN} --filetype=obj
|
||||
DEPENDS llvm-mc ${HIPRTC_GEN_PREPROCESSED} ${HIPRTC_GEN_MCIN})
|
||||
|
||||
# Create hiprtc-builtins library.
|
||||
add_library(hiprtc-builtins ${HIPRTC_GEN_OBJ})
|
||||
set_target_properties(hiprtc-builtins PROPERTIES
|
||||
CXX_STANDARD 14
|
||||
CXX_STANDARD_REQUIRED ON
|
||||
CXX_EXTENSIONS OFF
|
||||
POSITION_INDEPENDENT_CODE ON
|
||||
LIBRARY_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR}/lib
|
||||
LINKER_LANGUAGE CXX
|
||||
VERSION ${HIP_LIB_VERSION_STRING})
|
||||
|
||||
# Windows and Linux have different naming conventions.
|
||||
if(WIN32)
|
||||
# Windows uses DEF file to determine which symbols to expose.
|
||||
target_sources(hiprtc-builtins PRIVATE hiprtc-builtins.def)
|
||||
set_target_properties(hiprtc-builtins PROPERTIES
|
||||
OUTPUT_NAME "hiprtc-builtins64_${HIP_LIB_VERSION_MAJOR}${HIP_LIB_VERSION_MINOR}")
|
||||
# Since ${HIPRTC_GEN_OBJ} was manually generated with llvm-mc, /MT did not embed
|
||||
# libcmt.lib inside of the obj. So we need to manually set it as defaultlib.
|
||||
target_link_options(hiprtc-builtins PRIVATE "LINKER:/DEFAULTLIB:libcmt")
|
||||
else()
|
||||
# SOVERSION is only supported on Linux.
|
||||
set_target_properties(hiprtc-builtins PROPERTIES
|
||||
OUTPUT_NAME "hiprtc-builtins"
|
||||
SOVERSION ${HIP_LIB_VERSION_MAJOR})
|
||||
endif()
|
||||
|
||||
# Test the header file works with simple compilation.
|
||||
add_custom_command(
|
||||
OUTPUT ${HIPRTC_GEN_DIR}/tmp.bc
|
||||
COMMAND $<TARGET_FILE:clang> -O3 --rocm-path=${PROJECT_SOURCE_DIR}/include/.. -std=c++14 -nogpulib -nogpuinc -emit-llvm -c -isystem ${HIP_COMMON_INCLUDE_DIR} -isystem ${PROJECT_BINARY_DIR}/include -isystem ${CMAKE_CURRENT_SOURCE_DIR}/include --cuda-device-only -D__HIPCC_RTC__ --offload-arch=gfx906 -x hip-cpp-output ${HIPRTC_GEN_PREPROCESSED} -o ${HIPRTC_GEN_DIR}/tmp.bc
|
||||
DEPENDS clang ${HIPRTC_GEN_PREPROCESSED})
|
||||
|
||||
# FIXME: As a workaround, add hiprtc object into amdhip64, until we can
|
||||
# figure out how to link hiprtc-builtins into amdhip64. CMake approach is not working:
|
||||
# target_link_libraries(amdhip64 PUBLIC hiprtc-builtins)
|
||||
target_link_libraries(amdhip64 PRIVATE ${HIPRTC_GEN_OBJ})
|
||||
target_compile_definitions(amdhip64 PRIVATE __HIP_ENABLE_RTC)
|
||||
add_dependencies(amdhip64 hiprtc-builtins)
|
||||
# Windows will install the library to the bin directory instead.
|
||||
if(WIN32)
|
||||
install(TARGETS hiprtc-builtins DESTINATION ${CMAKE_INSTALL_BINDIR})
|
||||
else()
|
||||
install(TARGETS hiprtc-builtins DESTINATION lib)
|
||||
endif()
|
||||
endif()
|
||||
|
||||
#############################
|
||||
# Profiling API support
|
||||
#############################
|
||||
# Generate profiling API macros/structures header
|
||||
set(PROF_API_STR "${PROJECT_BINARY_DIR}/include/hip/amd_detail/hip_prof_str.h")
|
||||
set(PROF_API_HDR "${HIP_COMMON_INCLUDE_DIR}/hip/hip_runtime_api.h")
|
||||
set(PROF_API_SRC "${CMAKE_CURRENT_SOURCE_DIR}")
|
||||
set(PROF_API_GEN "${CMAKE_CURRENT_SOURCE_DIR}/hip_prof_gen.py")
|
||||
set(PROF_API_LOG "${PROJECT_BINARY_DIR}/hip_prof_gen.log.txt")
|
||||
|
||||
find_package(PythonInterp REQUIRED)
|
||||
add_custom_command(OUTPUT ${PROF_API_STR}
|
||||
COMMAND ${PYTHON_EXECUTABLE} ${PROF_API_GEN} -v -t --priv ${PROF_API_HDR} ${PROF_API_SRC} ${PROF_API_STR}
|
||||
DEPENDS ${PROF_API_HDR} ${PROF_API_GEN}
|
||||
COMMENT "Generating profiling primitives: ${PROF_API_STR}")
|
||||
|
||||
add_custom_target(gen-prof-api-str-header ALL
|
||||
DEPENDS ${PROF_API_STR}
|
||||
SOURCES ${PROF_API_HDR})
|
||||
|
||||
set_target_properties(amdhip64 PROPERTIES PUBLIC_HEADER ${PROF_API_STR})
|
||||
|
||||
option(USE_PROF_API ON "Enable roctracer integration")
|
||||
# Enable profiling API
|
||||
if(USE_PROF_API)
|
||||
find_path(PROF_API_HEADER_DIR prof_protocol.h
|
||||
HINTS
|
||||
${PROF_API_HEADER_PATH}
|
||||
PATHS
|
||||
${ROCM_PATH}/roctracer
|
||||
PATH_SUFFIXES
|
||||
include/ext)
|
||||
|
||||
if(NOT PROF_API_HEADER_DIR)
|
||||
message(WARNING "Profiling API header not found. Disabling roctracer integration. Use -DPROF_API_HEADER_PATH=<path to prof_protocol.h header>")
|
||||
else()
|
||||
target_compile_definitions(amdhip64 PUBLIC USE_PROF_API=1)
|
||||
target_include_directories(amdhip64 PUBLIC ${PROF_API_HEADER_DIR})
|
||||
message(STATUS "Profiling API: ${PROF_API_HEADER_DIR}")
|
||||
endif()
|
||||
endif()
|
||||
|
||||
add_dependencies(amdhip64 gen-prof-api-str-header)
|
||||
|
||||
add_custom_command(TARGET amdhip64 POST_BUILD COMMAND
|
||||
${CMAKE_COMMAND} -E copy ${PROJECT_BINARY_DIR}/.hipInfo ${PROJECT_BINARY_DIR}/lib/.hipInfo)
|
||||
add_custom_command(TARGET amdhip64 POST_BUILD COMMAND
|
||||
${CMAKE_COMMAND} -E copy_directory ${PROJECT_SOURCE_DIR}/include ${PROJECT_BINARY_DIR}/include)
|
||||
|
||||
add_library(host INTERFACE)
|
||||
target_link_libraries(host INTERFACE amdhip64)
|
||||
|
||||
add_library(device INTERFACE)
|
||||
target_link_libraries(device INTERFACE host)
|
||||
|
||||
# Current packaging assumes that HIP runtime will always be installed in /opt/rocm/lib
|
||||
# This is false to assume, because some distros like CentOS will use the lib64 directory instead of lib
|
||||
# Relying on CMake to choose the library directory for us will default in that case to lib64
|
||||
# Hence there will be a mismatch between where HIP is installed and where CMake thinks it is
|
||||
|
||||
INSTALL(TARGETS amdhip64 host device
|
||||
EXPORT hip-targets
|
||||
RUNTIME DESTINATION ${CMAKE_INSTALL_BINDIR}
|
||||
LIBRARY DESTINATION lib
|
||||
ARCHIVE DESTINATION ${CMAKE_INSTALL_LIBDIR}
|
||||
PUBLIC_HEADER DESTINATION ${CMAKE_INSTALL_INCLUDEDIR})
|
||||
INSTALL(EXPORT hip-targets DESTINATION ${CONFIG_PACKAGE_INSTALL_DIR} NAMESPACE hip::)
|
||||
|
||||
INSTALL(TARGETS amdhip64 host device
|
||||
EXPORT hip-lang-targets
|
||||
RUNTIME DESTINATION ${CMAKE_INSTALL_BINDIR}
|
||||
LIBRARY DESTINATION lib
|
||||
ARCHIVE DESTINATION ${CMAKE_INSTALL_LIBDIR}
|
||||
PUBLIC_HEADER DESTINATION ${CMAKE_INSTALL_INCLUDEDIR})
|
||||
INSTALL(EXPORT hip-lang-targets DESTINATION ${CONFIG_LANG_PACKAGE_INSTALL_DIR} NAMESPACE hip-lang::)
|
||||
|
||||
if(NOT WIN32)
|
||||
include(CMakePackageConfigHelpers)
|
||||
|
||||
configure_package_config_file(
|
||||
${HIP_COMMON_DIR}/hip-lang-config.cmake.in
|
||||
${CMAKE_CURRENT_BINARY_DIR}/hip-lang-config.cmake
|
||||
INSTALL_DESTINATION ${CONFIG_LANG_PACKAGE_INSTALL_DIR}
|
||||
PATH_VARS LIB_INSTALL_DIR INCLUDE_INSTALL_DIR BIN_INSTALL_DIR)
|
||||
|
||||
write_basic_package_version_file(
|
||||
${CMAKE_CURRENT_BINARY_DIR}/hip-lang-config-version.cmake
|
||||
VERSION "${HIP_VERSION_MAJOR}.${HIP_VERSION_MINOR}.${HIP_VERSION_GITDATE}"
|
||||
COMPATIBILITY SameMajorVersion)
|
||||
install(
|
||||
FILES
|
||||
${CMAKE_CURRENT_BINARY_DIR}/hip-lang-config.cmake
|
||||
${CMAKE_CURRENT_BINARY_DIR}/hip-lang-config-version.cmake
|
||||
DESTINATION
|
||||
${CONFIG_LANG_PACKAGE_INSTALL_DIR}/
|
||||
)
|
||||
endif()
|
||||
@@ -0,0 +1,51 @@
|
||||
# Copyright (c) 2020-2021 Advanced Micro Devices, Inc. All rights reserved.
|
||||
#
|
||||
# Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
# of this software and associated documentation files (the "Software"), to deal
|
||||
# in the Software without restriction, including without limitation the rights
|
||||
# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
# copies of the Software, and to permit persons to whom the Software is
|
||||
# furnished to do so, subject to the following conditions:
|
||||
#
|
||||
# The above copyright notice and this permission notice shall be included in
|
||||
# all copies or substantial portions of the Software.
|
||||
#
|
||||
# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
# THE SOFTWARE.
|
||||
|
||||
if(ROCCLR_FOUND)
|
||||
return()
|
||||
endif()
|
||||
|
||||
find_path(ROCCLR_INCLUDE_DIR top.hpp
|
||||
HINTS
|
||||
${ROCCLR_PATH}
|
||||
PATHS
|
||||
# gerrit repo name
|
||||
${CMAKE_SOURCE_DIR}/vdi
|
||||
${CMAKE_SOURCE_DIR}/../vdi
|
||||
${CMAKE_SOURCE_DIR}/../../vdi
|
||||
# github repo name
|
||||
${CMAKE_SOURCE_DIR}/ROCclr
|
||||
${CMAKE_SOURCE_DIR}/../ROCclr
|
||||
${CMAKE_SOURCE_DIR}/../../ROCclr
|
||||
# jenkins repo name
|
||||
${CMAKE_SOURCE_DIR}/rocclr
|
||||
${CMAKE_SOURCE_DIR}/../rocclr
|
||||
${CMAKE_SOURCE_DIR}/../../rocclr
|
||||
PATH_SUFFIXES
|
||||
include)
|
||||
|
||||
include(FindPackageHandleStandardArgs)
|
||||
find_package_handle_standard_args(ROCclr
|
||||
"\nROCclr not found"
|
||||
ROCCLR_INCLUDE_DIR)
|
||||
mark_as_advanced(ROCCLR_INCLUDE_DIR)
|
||||
|
||||
list(APPEND CMAKE_MODULE_PATH "${ROCCLR_INCLUDE_DIR}/../cmake")
|
||||
include(ROCclr)
|
||||
@@ -0,0 +1,86 @@
|
||||
# Copyright (C) 2021-2021 Advanced Micro Devices, Inc. All Rights Reserved.
|
||||
#
|
||||
# Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
# of this software and associated documentation files (the "Software"), to deal
|
||||
# in the Software without restriction, including without limitation the rights
|
||||
# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
# copies of the Software, and to permit persons to whom the Software is
|
||||
# furnished to do so, subject to the following conditions:
|
||||
#
|
||||
# The above copyright notice and this permission notice shall be included in
|
||||
# all copies or substantial portions of the Software.
|
||||
#
|
||||
# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
# THE SOFTWARE.
|
||||
|
||||
|
||||
###############################################################################
|
||||
# HIPRTC.cmake
|
||||
###############################################################################
|
||||
|
||||
# This file includes macros required to generate the hiprtc builtins library.
|
||||
|
||||
function(get_hiprtc_macros HIPRTC_DEFINES)
|
||||
set(${HIPRTC_DEFINES}
|
||||
"#define __device__ __attribute__((device))\n\
|
||||
#define __host__ __attribute__((host))\n\
|
||||
#define __global__ __attribute__((global))\n\
|
||||
#define __constant__ __attribute__((constant))\n\
|
||||
#define __shared__ __attribute__((shared))\n\
|
||||
|
||||
#define launch_bounds_impl0(requiredMaxThreadsPerBlock) \\\n\
|
||||
__attribute__((amdgpu_flat_work_group_size(1, requiredMaxThreadsPerBlock)))\n\
|
||||
#define launch_bounds_impl1(requiredMaxThreadsPerBlock, minBlocksPerMultiprocessor) \\\n\
|
||||
__attribute__((amdgpu_flat_work_group_size(1, requiredMaxThreadsPerBlock), \\\n\
|
||||
amdgpu_waves_per_eu(minBlocksPerMultiprocessor)))\n\
|
||||
#define select_impl_(_1, _2, impl_, ...) impl_\n\
|
||||
#define __launch_bounds__(...) \\\n\
|
||||
select_impl_(__VA_ARGS__, launch_bounds_impl1, launch_bounds_impl0)(__VA_ARGS__)"
|
||||
PARENT_SCOPE)
|
||||
endfunction(get_hiprtc_macros)
|
||||
|
||||
# To allow concatenating above macros during build time, call this file in script mode.
|
||||
if(HIPRTC_ADD_MACROS)
|
||||
message(STATUS "Appending hiprtc macros to ${HIPRTC_PREPROCESSED_FILE}.")
|
||||
get_hiprtc_macros(HIPRTC_DEFINES)
|
||||
FILE(APPEND ${HIPRTC_PREPROCESSED_FILE} "${HIPRTC_DEFINES}")
|
||||
endif()
|
||||
|
||||
macro(generate_hiprtc_header HiprtcHeader)
|
||||
FILE(WRITE ${HiprtcHeader}
|
||||
"#pragma push_macro(\"CHAR_BIT\")\n\
|
||||
#pragma push_macro(\"INT_MAX\")\n\
|
||||
#define CHAR_BIT __CHAR_BIT__\n\
|
||||
#define INT_MAX __INTMAX_MAX__\n\
|
||||
#include \"hip/hip_runtime.h\"\n\
|
||||
#include \"hip/hip_fp16.h\"\n\
|
||||
#pragma pop_macro(\"CHAR_BIT\")\n\
|
||||
#pragma pop_macro(\"INT_MAX\")")
|
||||
endmacro(generate_hiprtc_header)
|
||||
|
||||
macro(generate_hiprtc_mcin HiprtcMcin HiprtcPreprocessedInput)
|
||||
if(WIN32)
|
||||
set(HIPRTC_TYPE_LINUX_ONLY "")
|
||||
else()
|
||||
set(HIPRTC_TYPE_LINUX_ONLY
|
||||
" .type __hipRTC_header,@object\n"
|
||||
" .type __hipRTC_header_size,@object")
|
||||
endif()
|
||||
FILE(WRITE ${HiprtcMcin}
|
||||
"// Automatically generated script for HIPRTC.\n\
|
||||
${HIPRTC_TYPE_LINUX_ONLY}\n\
|
||||
.section .hipRTC_header,\"a\"\n\
|
||||
.globl __hipRTC_header\n\
|
||||
.globl __hipRTC_header_size\n\
|
||||
.p2align 3\n\
|
||||
__hipRTC_header:\n\
|
||||
.incbin \"${HiprtcPreprocessedInput}\"\n\
|
||||
__hipRTC_header_size:\n\
|
||||
.long __hipRTC_header_size - __hipRTC_header\n")
|
||||
endmacro(generate_hiprtc_mcin)
|
||||
|
||||
Reference in New Issue
Block a user