SWDEV-286884 - Add CMake hiprtc-builtins library generation part 1
Add the CMake support for generating the hiprtc-builtins library. Replaces the previous execution of shell script. For backwards compatibility, link hiprtc pre-processed object to amdhip64, and support versioning in name. Add libcmt.lib to Windows link step, due to manually generated .obj from llvm-mc. Change-Id: I267be3cf4b241840b35f7f27a0b8659530108b0e
This commit is contained in:
Aaron En Ye Shi
+69
-17
@@ -171,28 +171,80 @@ endif()
|
||||
|
||||
# Enable preprocessed hiprtc-builtins library
|
||||
if(__HIP_ENABLE_RTC)
|
||||
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}/../../..")
|
||||
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 ${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 src/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 ${PROJECT_SOURCE_DIR}/include -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)
|
||||
set(HIPRTC_LIB_NAME "hiprtc-builtins64_${HIP_LIB_VERSION_MAJOR}${HIP_LIB_VERSION_MINOR}.dll")
|
||||
install(TARGETS hiprtc-builtins DESTINATION ${CMAKE_INSTALL_BINDIR})
|
||||
else()
|
||||
set(HIPRTC_LIB_NAME "libhiprtc-builtins.so.${HIP_LIB_VERSION_MAJOR}.${HIP_LIB_VERSION_MINOR}")
|
||||
install(TARGETS hiprtc-builtins DESTINATION lib)
|
||||
endif()
|
||||
execute_process(
|
||||
COMMAND sh -c "mkdir -p ${PROJECT_BINARY_DIR}/lib; ${HIP_COMMON_DIR}/bin/hip_embed_pch.sh ${PROJECT_BINARY_DIR}/include ${PROJECT_SOURCE_DIR}/include ${CMAKE_CURRENT_SOURCE_DIR}/include ${HIP_LLVM_ROOT} -r ${PROJECT_BINARY_DIR}/lib/${HIPRTC_LIB_NAME}"
|
||||
COMMAND_ECHO STDERR
|
||||
RESULT_VARIABLE EMBED_RTC_RC
|
||||
)
|
||||
if (EMBED_RTC_RC AND NOT EMBED_RTC_RC EQUAL 0)
|
||||
message(FATAL_ERROR "Failed to create hiprtc shared lib")
|
||||
endif()
|
||||
install(FILES ${PROJECT_BINARY_DIR}/lib/${HIPRTC_LIB_NAME} DESTINATION lib)
|
||||
endif()
|
||||
|
||||
#############################
|
||||
|
||||
@@ -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)
|
||||
|
||||
@@ -0,0 +1,4 @@
|
||||
EXPORTS
|
||||
__hipRTC_header
|
||||
__hipRTC_header_size
|
||||
|
||||
Reference in New Issue
Block a user