diff --git a/projects/clr/hipamd/CMakeLists.txt b/projects/clr/hipamd/CMakeLists.txt index e66666519f..76393736e2 100755 --- a/projects/clr/hipamd/CMakeLists.txt +++ b/projects/clr/hipamd/CMakeLists.txt @@ -47,7 +47,6 @@ set(CMAKE_INSTALL_RPATH_USE_LINK_PATH TRUE) ############################# option(BUILD_HIPIFY_CLANG "Enable building the CUDA->HIP converter" OFF) option(__HIP_ENABLE_PCH "Enable/Disable pre-compiled hip headers" ON) -option(__HIP_ENABLE_RTC "Enable/Disable pre-processed hiprtc shared lib" ON) option(__HIP_USE_CMPXCHG_FOR_FP_ATOMICS "Enable/Disable FP Atomics version" OFF) option(HIP_OFFICIAL_BUILD "Enable/Disable for mainline/staging builds" OFF) set(HIPCC_BIN_DIR "" CACHE STRING "HIPCC and HIPCONFIG binary directories") diff --git a/projects/clr/hipamd/packaging/CMakeLists.txt b/projects/clr/hipamd/packaging/CMakeLists.txt index 58ba334104..9fe824cf0b 100644 --- a/projects/clr/hipamd/packaging/CMakeLists.txt +++ b/projects/clr/hipamd/packaging/CMakeLists.txt @@ -44,11 +44,12 @@ if(HIP_PLATFORM STREQUAL "amd" ) install(FILES ${CMAKE_BINARY_DIR}/lib/libamdhip64.so DESTINATION lib COMPONENT binary) install(FILES ${CMAKE_BINARY_DIR}/lib/libamdhip64.so.${HIP_LIB_VERSION_MAJOR} DESTINATION lib COMPONENT binary) install(FILES ${CMAKE_BINARY_DIR}/lib/libamdhip64.so.${HIP_LIB_VERSION_STRING} DESTINATION lib COMPONENT binary) - if(__HIP_ENABLE_RTC) - install(FILES ${CMAKE_BINARY_DIR}/lib/libhiprtc-builtins.so DESTINATION lib COMPONENT binary) - install(FILES ${CMAKE_BINARY_DIR}/lib/libhiprtc-builtins.so.${HIP_LIB_VERSION_MAJOR} DESTINATION lib COMPONENT binary) - install(FILES ${CMAKE_BINARY_DIR}/lib/libhiprtc-builtins.so.${HIP_LIB_VERSION_STRING} DESTINATION lib COMPONENT binary) - endif() + install(FILES ${CMAKE_BINARY_DIR}/lib/libhiprtc.so DESTINATION lib COMPONENT binary) + install(FILES ${CMAKE_BINARY_DIR}/lib/libhiprtc.so.${HIP_LIB_VERSION_MAJOR} DESTINATION lib COMPONENT binary) + install(FILES ${CMAKE_BINARY_DIR}/lib/libhiprtc.so.${HIP_LIB_VERSION_STRING} DESTINATION lib COMPONENT binary) + install(FILES ${CMAKE_BINARY_DIR}/lib/libhiprtc-builtins.so DESTINATION lib COMPONENT binary) + install(FILES ${CMAKE_BINARY_DIR}/lib/libhiprtc-builtins.so.${HIP_LIB_VERSION_MAJOR} DESTINATION lib COMPONENT binary) + install(FILES ${CMAKE_BINARY_DIR}/lib/libhiprtc-builtins.so.${HIP_LIB_VERSION_STRING} DESTINATION lib COMPONENT binary) else() install(FILES ${CMAKE_BINARY_DIR}/lib/libamdhip64.a DESTINATION lib COMPONENT binary) endif()#End BUILD_SHARED_LIBS diff --git a/projects/clr/hipamd/src/CMakeLists.txt b/projects/clr/hipamd/src/CMakeLists.txt index ac88ddf5d4..c747b365ec 100644 --- a/projects/clr/hipamd/src/CMakeLists.txt +++ b/projects/clr/hipamd/src/CMakeLists.txt @@ -1,4 +1,4 @@ -# Copyright (c) 2020 - 2021 Advanced Micro Devices, Inc. All rights reserved. +# Copyright (c) 2020 - 2022 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 @@ -110,7 +110,6 @@ target_sources(amdhip64 PRIVATE hip_peer.cpp hip_platform.cpp hip_profile.cpp - hip_rtc.cpp hip_stream_ops.cpp hip_stream.cpp hip_surface.cpp @@ -147,10 +146,6 @@ target_include_directories(amdhip64 target_compile_definitions(amdhip64 PRIVATE __HIP_PLATFORM_AMD__) target_link_libraries(amdhip64 PRIVATE ${OPENGL_LIBRARIES}) 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. @@ -185,80 +180,20 @@ if(__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 ${ROCM_PATH}/llvm) - find_package(Clang REQUIRED CONFIG PATHS ${ROCM_PATH}/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}") +set(HIPRTC_OBJECTS) +# Add hiprtc +add_subdirectory(hiprtc) - # 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 $ -O3 --rocm-path=${PROJECT_SOURCE_DIR}/include/.. -std=c++17 -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 $ -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}) +if(NOT WIN32) + if(BUILD_SHARED_LIBS) + target_link_libraries(amdhip64 PRIVATE ${HIPRTC_OBJECTS}) + target_compile_definitions(amdhip64 PRIVATE __HIP_ENABLE_RTC) + add_dependencies(amdhip64 hiprtc-builtins) + INSTALL(TARGETS hiprtc-builtins + RUNTIME DESTINATION ${CMAKE_INSTALL_BINDIR} + LIBRARY DESTINATION lib + ARCHIVE DESTINATION ${CMAKE_INSTALL_LIBDIR}) endif() - - # Test the header file works with simple compilation. - add_custom_command( - OUTPUT ${HIPRTC_GEN_DIR}/tmp.bc - COMMAND $ -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) - install(TARGETS hiprtc-builtins - RUNTIME DESTINATION ${CMAKE_INSTALL_BINDIR} - LIBRARY DESTINATION lib - ARCHIVE DESTINATION ${CMAKE_INSTALL_LIBDIR}) endif() ############################# diff --git a/projects/clr/hipamd/src/amdhip.def b/projects/clr/hipamd/src/amdhip.def index 152296fe06..c5ed8e6a12 100644 --- a/projects/clr/hipamd/src/amdhip.def +++ b/projects/clr/hipamd/src/amdhip.def @@ -239,16 +239,6 @@ hipDestroySurfaceObject hipInitActivityCallback hipEnableActivityCallback hipGetCmdName -hiprtcAddNameExpression -hiprtcCompileProgram -hiprtcCreateProgram -hiprtcDestroyProgram -hiprtcGetLoweredName -hiprtcGetProgramLog -hiprtcGetProgramLogSize -hiprtcGetCode -hiprtcGetCodeSize -hiprtcGetErrorString hipMipmappedArrayCreate hipMallocMipmappedArray hipMipmappedArrayDestroy diff --git a/projects/clr/hipamd/src/hiprtc/CMakeLists.txt b/projects/clr/hipamd/src/hiprtc/CMakeLists.txt new file mode 100644 index 0000000000..f72758ccd9 --- /dev/null +++ b/projects/clr/hipamd/src/hiprtc/CMakeLists.txt @@ -0,0 +1,195 @@ +# Copyright (c) 2020 - 2022 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. + +# This project builds hiprtc +# If ever this is to be a different lib living in different folder +# Please read this part +# Depends on: rocclr, so find_package(rocclr) will be required +# Building hip header requires hip include folders with hip_version.h + +cmake_minimum_required(VERSION 3.16.1) +option(BUILD_SHARED_LIBS "Build the shared library" ON) + +list(APPEND CMAKE_MODULE_PATH "${CMAKE_CURRENT_LIST_DIR}/cmake") + +if(BUILD_SHARED_LIBS) + add_library(hiprtc 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 hiprtc POST_BUILD COMMAND ${CMAKE_STRIP} $) + endif() +else() + add_library(hiprtc STATIC $) +endif() + +set_target_properties(hiprtc PROPERTIES + CXX_STANDARD 17 + CXX_STANDARD_REQUIRED ON + CXX_EXTENSIONS OFF + POSITION_INDEPENDENT_CODE ON + LIBRARY_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR}/lib + ARCHIVE_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR}/lib) + +if(NOT WIN32) + if(BUILD_SHARED_LIBS) + set_target_properties(hiprtc PROPERTIES + VERSION ${HIP_LIB_VERSION_STRING} + SOVERSION ${HIP_LIB_VERSION_MAJOR}) + endif() +endif() + +# Create HIPRTC object library +if(BUILD_SHARED_LIBS) + add_library(hiprtcobject OBJECT hiprtc.cpp hiprtcComgrHelper.cpp hiprtcInternal.cpp) +endif() + +set_target_properties(hiprtcobject PROPERTIES + CXX_STANDARD 17 + CXX_STANDARD_REQUIRED ON + CXX_EXTENSIONS OFF + POSITION_INDEPENDENT_CODE ON) + +target_include_directories(hiprtcobject + PRIVATE + ${HIP_COMMON_INCLUDE_DIR} + ${PROJECT_SOURCE_DIR}/include + ${PROJECT_BINARY_DIR}/include) + +if(BUILD_SHARED_LIBS) + if(WIN32) + target_sources(hiprtc PRIVATE hiprtc.def) + else() + target_link_libraries(hiprtcobject PRIVATE "-Wl,--version-script=${CMAKE_CURRENT_LIST_DIR}/hiprtc.map.in") + set_target_properties(hiprtcobject PROPERTIES LINK_DEPENDS "${CMAKE_CURRENT_LIST_DIR}/hiprtc.map.in") + endif() +endif() + +if(WIN32) + target_link_libraries(hiprtc PRIVATE Dbghelp.lib) +endif() + +target_link_libraries(hiprtcobject PRIVATE ${CMAKE_DL_LIBS}) + +if(BUILD_SHARED_LIBS) + target_link_libraries(hiprtcobject PRIVATE rocclr) +else() + target_compile_definitions(hiprtcobject PRIVATE $) + target_include_directories(hiprtcobject PRIVATE $) +endif() + +target_compile_definitions(hiprtcobject PRIVATE __HIP_PLATFORM_AMD__) + +add_to_config(_versionInfo HIP_PACKAGING_VERSION_PATCH) +add_to_config(_versionInfo CPACK_DEBIAN_PACKAGE_RELEASE) +add_to_config(_versionInfo CPACK_RPM_PACKAGE_RELEASE) + +add_to_config(_versionInfo HIP_VERSION_MAJOR) +add_to_config(_versionInfo HIP_VERSION_MINOR) +add_to_config(_versionInfo HIP_VERSION_PATCH) +add_to_config(_versionInfo HIP_VERSION_GITHASH) + +# Enable preprocessed hiprtc-builtins library +include(HIPRTC RESULT_VARIABLE HIPRTC_CMAKE) +# Requires clang and llvm-mc to create this library. +find_package(LLVM REQUIRED CONFIG PATHS ${ROCM_PATH}/llvm) +find_package(Clang REQUIRED CONFIG PATHS ${ROCM_PATH}/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 $ -O3 --rocm-path=${PROJECT_SOURCE_DIR}/include/.. -std=c++17 -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 $ -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 $ -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}) + +target_link_libraries(hiprtcobject PRIVATE ${HIPRTC_GEN_OBJ}) +target_compile_definitions(hiprtcobject PRIVATE __HIP_ENABLE_RTC) +target_link_libraries(hiprtc PRIVATE hiprtcobject) + +# As a temporary workaround adding hiprtc sources to amdhip64 using target_sources, to avoid +# jenkins failure. Once, jenkins upgrades to 3.21 or higher, hiprtcobject can be appended to +# HIPRTC_OBJECTS below which links to amdhip64 +if(NOT WIN32) + target_sources(amdhip64 PRIVATE hiprtc.cpp hiprtcComgrHelper.cpp hiprtcInternal.cpp) +endif() + +list(APPEND HIPRTC_OBJECTS ${HIPRTC_GEN_OBJ}) +set(HIPRTC_OBJECTS ${HIPRTC_OBJECTS} PARENT_SCOPE) + +add_dependencies(hiprtc hiprtc-builtins) +install(TARGETS hiprtc-builtins + RUNTIME DESTINATION ${CMAKE_INSTALL_BINDIR} + LIBRARY DESTINATION lib + ARCHIVE DESTINATION ${CMAKE_INSTALL_LIBDIR}) + +INSTALL(TARGETS hiprtc + EXPORT hiprtc-targets + RUNTIME DESTINATION ${CMAKE_INSTALL_BINDIR} + LIBRARY DESTINATION lib + ARCHIVE DESTINATION ${CMAKE_INSTALL_LIBDIR} + PUBLIC_HEADER DESTINATION ${CMAKE_INSTALL_INCLUDEDIR}) diff --git a/projects/clr/hipamd/src/cmake/HIPRTC.cmake b/projects/clr/hipamd/src/hiprtc/cmake/HIPRTC.cmake similarity index 95% rename from projects/clr/hipamd/src/cmake/HIPRTC.cmake rename to projects/clr/hipamd/src/hiprtc/cmake/HIPRTC.cmake index bc014360c1..0443588a4b 100644 --- a/projects/clr/hipamd/src/cmake/HIPRTC.cmake +++ b/projects/clr/hipamd/src/hiprtc/cmake/HIPRTC.cmake @@ -1,4 +1,4 @@ -# Copyright (c) 2021 - 2021 Advanced Micro Devices, Inc. All Rights Reserved. +# Copyright (c) 2021 - 2022 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 @@ -35,6 +35,8 @@ function(get_hiprtc_macros HIPRTC_DEFINES) #define __constant__ __attribute__((constant))\n\ #define __shared__ __attribute__((shared))\n\ #define __align__(x) __attribute__((aligned(x)))\n\ +#define __noinline__ __attribute__((noinline))\n\ +#define __forceinline__ inline __attribute__((always_inline))\n\ #define launch_bounds_impl0(requiredMaxThreadsPerBlock) \\\n\ __attribute__((amdgpu_flat_work_group_size(1, requiredMaxThreadsPerBlock)))\n\ diff --git a/projects/clr/hipamd/src/hiprtc-builtins.def b/projects/clr/hipamd/src/hiprtc/hiprtc-builtins.def similarity index 100% rename from projects/clr/hipamd/src/hiprtc-builtins.def rename to projects/clr/hipamd/src/hiprtc/hiprtc-builtins.def diff --git a/projects/clr/hipamd/src/hiprtc/hiprtc.cpp b/projects/clr/hipamd/src/hiprtc/hiprtc.cpp new file mode 100644 index 0000000000..c1c17c38bb --- /dev/null +++ b/projects/clr/hipamd/src/hiprtc/hiprtc.cpp @@ -0,0 +1,226 @@ +/* +Copyright (c) 2022 - Present 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. +*/ + +#include +#include "hiprtcInternal.hpp" + +namespace hiprtc { +thread_local hiprtcResult g_lastRtcError = HIPRTC_SUCCESS; +} + +const char* hiprtcGetErrorString(hiprtcResult x) { + switch (x) { + case HIPRTC_SUCCESS: + return "HIPRTC_SUCCESS"; + case HIPRTC_ERROR_OUT_OF_MEMORY: + return "HIPRTC_ERROR_OUT_OF_MEMORY"; + case HIPRTC_ERROR_PROGRAM_CREATION_FAILURE: + return "HIPRTC_ERROR_PROGRAM_CREATION_FAILURE"; + case HIPRTC_ERROR_INVALID_INPUT: + return "HIPRTC_ERROR_INVALID_INPUT"; + case HIPRTC_ERROR_INVALID_PROGRAM: + return "HIPRTC_ERROR_INVALID_PROGRAM"; + case HIPRTC_ERROR_INVALID_OPTION: + return "HIPRTC_ERROR_INVALID_OPTION"; + case HIPRTC_ERROR_COMPILATION: + return "HIPRTC_ERROR_COMPILATION"; + case HIPRTC_ERROR_BUILTIN_OPERATION_FAILURE: + return "HIPRTC_ERROR_BUILTIN_OPERATION_FAILURE"; + case HIPRTC_ERROR_NO_NAME_EXPRESSIONS_AFTER_COMPILATION: + return "HIPRTC_ERROR_NO_NAME_EXPRESSIONS_AFTER_COMPILATION"; + case HIPRTC_ERROR_NO_LOWERED_NAMES_BEFORE_COMPILATION: + return "HIPRTC_ERROR_NO_LOWERED_NAMES_BEFORE_COMPILATION"; + case HIPRTC_ERROR_NAME_EXPRESSION_NOT_VALID: + return "HIPRTC_ERROR_NAME_EXPRESSION_NOT_VALID"; + case HIPRTC_ERROR_INTERNAL_ERROR: + return "HIPRTC_ERROR_INTERNAL_ERROR"; + default: + LogPrintfError("Invalid HIPRTC error code: %d \n", x); + return nullptr; + }; + + return nullptr; +} + + +hiprtcResult hiprtcCreateProgram(hiprtcProgram* prog, const char* src, const char* name, + int numHeaders, const char** headers, const char** headerNames) { + HIPRTC_INIT_API(prog, src, name, numHeaders, headers, headerNames); + + if (prog == nullptr) { + HIPRTC_RETURN(HIPRTC_ERROR_INVALID_PROGRAM); + } + if (numHeaders < 0) { + HIPRTC_RETURN(HIPRTC_ERROR_INVALID_INPUT); + } + if (numHeaders && (headers == nullptr || headerNames == nullptr)) { + HIPRTC_RETURN(HIPRTC_ERROR_INVALID_INPUT); + } + if(name == nullptr) { + HIPRTC_RETURN(HIPRTC_ERROR_INVALID_INPUT); + } + + auto* rtcProgram = new hiprtc::RTCProgram(std::string(name)); + if (rtcProgram == nullptr) { + HIPRTC_RETURN(HIPRTC_ERROR_PROGRAM_CREATION_FAILURE); + } + + if (!rtcProgram->addSource(std::string(src), std::string("CompileSource"))) { + delete rtcProgram; + HIPRTC_RETURN(HIPRTC_ERROR_INVALID_INPUT); + } + + for (int i = 0; i < numHeaders; i++) { + if (!rtcProgram->addHeader(std::string(headers[i]), std::string(headerNames[i]))) { + delete rtcProgram; + HIPRTC_RETURN(HIPRTC_ERROR_INVALID_INPUT); + } + } + + *prog = hiprtc::RTCProgram::as_hiprtcProgram(rtcProgram); + + HIPRTC_RETURN(HIPRTC_SUCCESS); +} + +hiprtcResult hiprtcCompileProgram(hiprtcProgram prog, int numOptions, const char** options) { + HIPRTC_INIT_API(prog, numOptions, options); + + auto* rtcProgram = hiprtc::RTCProgram::as_RTCProgram(prog); + + std::vector opt; + opt.reserve(numOptions); + for (int i = 0; i < numOptions; i++) { + opt.push_back(std::string(options[i])); + } + + if (!rtcProgram->compile(opt)) { + HIPRTC_RETURN(HIPRTC_ERROR_COMPILATION); + } + + HIPRTC_RETURN(HIPRTC_SUCCESS); +} + +hiprtcResult hiprtcAddNameExpression(hiprtcProgram prog, const char* name_expression) { + HIPRTC_INIT_API(prog, name_expression); + + if (name_expression == nullptr) { + HIPRTC_RETURN(HIPRTC_ERROR_INVALID_INPUT); + } + auto* rtcProgram = hiprtc::RTCProgram::as_RTCProgram(prog); + std::string name = name_expression; + if (!rtcProgram->trackMangledName(name)) { + HIPRTC_RETURN(HIPRTC_ERROR_INVALID_INPUT); + } + + HIPRTC_RETURN(HIPRTC_SUCCESS); +} + +hiprtcResult hiprtcGetLoweredName(hiprtcProgram prog, const char* name_expression, + const char** loweredName) { + HIPRTC_INIT_API(prog, name_expression, loweredName); + + if (name_expression == nullptr || loweredName == nullptr) { + HIPRTC_RETURN(HIPRTC_ERROR_INVALID_INPUT); + } + + auto* rtcProgram = hiprtc::RTCProgram::as_RTCProgram(prog); + + if (!rtcProgram->getDemangledName(name_expression, loweredName)) { + return HIPRTC_RETURN(HIPRTC_ERROR_NAME_EXPRESSION_NOT_VALID); + } + + HIPRTC_RETURN(HIPRTC_SUCCESS); +} + +hiprtcResult hiprtcDestroyProgram(hiprtcProgram* prog) { + HIPRTC_INIT_API(prog); + if (prog == nullptr) { + HIPRTC_RETURN(HIPRTC_ERROR_INVALID_INPUT); + } + auto* rtcProgram = hiprtc::RTCProgram::as_RTCProgram(*prog); + delete rtcProgram; + HIPRTC_RETURN(HIPRTC_SUCCESS); +} + +hiprtcResult hiprtcGetCodeSize(hiprtcProgram prog, size_t* binarySizeRet) { + HIPRTC_INIT_API(prog, binarySizeRet); + + if (binarySizeRet == nullptr) { + HIPRTC_RETURN(HIPRTC_ERROR_INVALID_INPUT); + } + auto* rtcProgram = hiprtc::RTCProgram::as_RTCProgram(prog); + + *binarySizeRet = rtcProgram->getExecSize(); + + HIPRTC_RETURN(HIPRTC_SUCCESS); +} + +hiprtcResult hiprtcGetCode(hiprtcProgram prog, char* binaryMem) { + HIPRTC_INIT_API(prog, binaryMem); + + if (binaryMem == nullptr) { + HIPRTC_RETURN(HIPRTC_ERROR_INVALID_INPUT); + } + auto* rtcProgram = hiprtc::RTCProgram::as_RTCProgram(prog); + auto binary = rtcProgram->getExec(); + ::memcpy(binaryMem, binary.data(), binary.size()); + + HIPRTC_RETURN(HIPRTC_SUCCESS); +} + +hiprtcResult hiprtcGetProgramLog(hiprtcProgram prog, char* dst) { + HIPRTC_INIT_API(prog, dst); + if (dst == nullptr) { + HIPRTC_RETURN(HIPRTC_ERROR_INVALID_INPUT); + } + auto* rtcProgram = hiprtc::RTCProgram::as_RTCProgram(prog); + auto log = rtcProgram->getLog(); + ::memcpy(dst, log.data(), log.size()); + + HIPRTC_RETURN(HIPRTC_SUCCESS); +} + +hiprtcResult hiprtcGetProgramLogSize(hiprtcProgram prog, size_t* logSizeRet) { + HIPRTC_INIT_API(prog, logSizeRet); + if (logSizeRet == nullptr) { + HIPRTC_RETURN(HIPRTC_ERROR_INVALID_INPUT); + } + auto* rtcProgram = hiprtc::RTCProgram::as_RTCProgram(prog); + + *logSizeRet = rtcProgram->getLogSize(); + + HIPRTC_RETURN(HIPRTC_SUCCESS); +} + +hiprtcResult hiprtcVersion(int* major, int* minor) { + HIPRTC_INIT_API(major, minor); + + if (major == nullptr || minor == nullptr) { + HIPRTC_RETURN(HIPRTC_ERROR_INVALID_INPUT); + } + + // TODO add actual version, what do these numbers mean + *major = 9; + *minor = 0; + + HIPRTC_RETURN(HIPRTC_SUCCESS); +} diff --git a/projects/clr/hipamd/src/hiprtc/hiprtc.def b/projects/clr/hipamd/src/hiprtc/hiprtc.def new file mode 100644 index 0000000000..f69854720b --- /dev/null +++ b/projects/clr/hipamd/src/hiprtc/hiprtc.def @@ -0,0 +1,11 @@ +EXPORTS +hiprtcAddNameExpression +hiprtcCompileProgram +hiprtcCreateProgram +hiprtcDestroyProgram +hiprtcGetLoweredName +hiprtcGetProgramLog +hiprtcGetProgramLogSize +hiprtcGetCode +hiprtcGetCodeSize +hiprtcGetErrorString \ No newline at end of file diff --git a/projects/clr/hipamd/src/hiprtc/hiprtc.map.in b/projects/clr/hipamd/src/hiprtc/hiprtc.map.in new file mode 100644 index 0000000000..abb3fe8d1f --- /dev/null +++ b/projects/clr/hipamd/src/hiprtc/hiprtc.map.in @@ -0,0 +1,16 @@ +{ +global: + hiprtcCompileProgram; + hiprtcCreateProgram; + hiprtcDestroyProgram; + hiprtcGetLoweredName; + hiprtcGetProgramLog; + hiprtcGetProgramLogSize; + hiprtcGetCode; + hiprtcGetCodeSize; + hiprtcGetErrorString; + hiprtcAddNameExpression; + hiprtcVersion; +local: + *; +}; diff --git a/projects/clr/hipamd/src/hiprtc/hiprtcComgrHelper.cpp b/projects/clr/hipamd/src/hiprtc/hiprtcComgrHelper.cpp new file mode 100644 index 0000000000..88af22a32b --- /dev/null +++ b/projects/clr/hipamd/src/hiprtc/hiprtcComgrHelper.cpp @@ -0,0 +1,538 @@ +/* +Copyright (c) 2022 - Present 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. +*/ + +#include "hiprtcComgrHelper.hpp" + +namespace hiprtc { + +namespace helpers { +bool addCodeObjData(amd_comgr_data_set_t& input, const std::vector& source, + const std::string& name, const amd_comgr_data_kind_t type) { + amd_comgr_data_t data; + + if (auto res = amd::Comgr::create_data(type, &data); res != AMD_COMGR_STATUS_SUCCESS) { + return false; + } + + if (auto res = amd::Comgr::set_data(data, source.size(), source.data()); + res != AMD_COMGR_STATUS_SUCCESS) { + amd::Comgr::release_data(data); + return false; + } + + if (auto res = amd::Comgr::set_data_name(data, name.c_str()); res != AMD_COMGR_STATUS_SUCCESS) { + amd::Comgr::release_data(data); + return false; + } + + if (auto res = amd::Comgr::data_set_add(input, data); res != AMD_COMGR_STATUS_SUCCESS) { + amd::Comgr::release_data(data); + return false; + } + amd::Comgr::release_data(data); // Release from our end after setting the input + + return true; +} + +bool extractBuildLog(amd_comgr_data_set_t dataSet, std::string& buildLog) { + size_t count; + if (auto res = amd::Comgr::action_data_count(dataSet, AMD_COMGR_DATA_KIND_LOG, &count); + res != AMD_COMGR_STATUS_SUCCESS) { + return false; + } + + std::vector log; + if (count > 0) { + if (!extractByteCodeBinary(dataSet, AMD_COMGR_DATA_KIND_LOG, log)) return false; + buildLog.insert(buildLog.end(), log.data(), log.data() + log.size()); + } + return true; +} + +bool extractByteCodeBinary(const amd_comgr_data_set_t inDataSet, + const amd_comgr_data_kind_t dataKind, std::vector& bin) { + amd_comgr_data_t binaryData; + + if (auto res = amd::Comgr::action_data_get_data(inDataSet, dataKind, 0, &binaryData); + res != AMD_COMGR_STATUS_SUCCESS) { + return false; + } + + size_t binarySize = 0; + if (auto res = amd::Comgr::get_data(binaryData, &binarySize, NULL); + res != AMD_COMGR_STATUS_SUCCESS) { + amd::Comgr::release_data(binaryData); + return false; + } + + size_t bufSize = (dataKind == AMD_COMGR_DATA_KIND_LOG) ? binarySize + 1 : binarySize; + + char* binary = new char[bufSize]; + if (binary == nullptr) { + amd::Comgr::release_data(binaryData); + return false; + } + + + if (auto res = amd::Comgr::get_data(binaryData, &binarySize, binary); + res != AMD_COMGR_STATUS_SUCCESS) { + delete[] binary; + amd::Comgr::release_data(binaryData); + return false; + } + + if (dataKind == AMD_COMGR_DATA_KIND_LOG) { + binary[binarySize] = '\0'; + } + + amd::Comgr::release_data(binaryData); + + bin.reserve(binarySize); + bin.assign(binary, binary + binarySize); + delete[] binary; + + return true; +} + +bool createAction(amd_comgr_action_info_t& action, std::vector& options, + const std::string& isa, const amd_comgr_language_t lang) { + if (auto res = amd::Comgr::create_action_info(&action); res != AMD_COMGR_STATUS_SUCCESS) { + return false; + } + + if (lang != AMD_COMGR_LANGUAGE_NONE) { + if (auto res = amd::Comgr::action_info_set_language(action, lang); + res != AMD_COMGR_STATUS_SUCCESS) { + amd::Comgr::destroy_action_info(action); + return false; + } + } + + if (auto res = amd::Comgr::action_info_set_isa_name(action, isa.c_str()); + res != AMD_COMGR_STATUS_SUCCESS) { + amd::Comgr::destroy_action_info(action); + return false; + } + + std::vector optionsArgv; + optionsArgv.reserve(options.size()); + for (auto& option : options) { + optionsArgv.push_back(option.c_str()); + } + + if (auto res = + amd::Comgr::action_info_set_option_list(action, optionsArgv.data(), optionsArgv.size()); + res != AMD_COMGR_STATUS_SUCCESS) { + amd::Comgr::destroy_action_info(action); + return res; + } + + if (auto res = amd::Comgr::action_info_set_logging(action, true); + res != AMD_COMGR_STATUS_SUCCESS) { + amd::Comgr::destroy_action_info(action); + return res; + } + + return AMD_COMGR_STATUS_SUCCESS; +} + +bool compileToBitCode(const amd_comgr_data_set_t compileInputs, const std::string& isa, + std::vector& compileOptions, std::string& buildLog, + std::vector& LLVMBitcode) { + amd_comgr_language_t lang = AMD_COMGR_LANGUAGE_HIP; + amd_comgr_action_info_t action; + amd_comgr_data_set_t output; + amd_comgr_data_set_t input = compileInputs; + + if (auto res = createAction(action, compileOptions, isa, lang); res != AMD_COMGR_STATUS_SUCCESS) { + return false; + } + + if (auto res = amd::Comgr::create_data_set(&output); res != AMD_COMGR_STATUS_SUCCESS) { + amd::Comgr::destroy_action_info(action); + return false; + } + + if (auto res = + amd::Comgr::do_action(AMD_COMGR_ACTION_COMPILE_SOURCE_TO_BC, action, input, output); + res != AMD_COMGR_STATUS_SUCCESS) { + extractBuildLog(output, buildLog); + amd::Comgr::destroy_action_info(action); + amd::Comgr::destroy_data_set(output); + return false; + } + + if (!extractBuildLog(output, buildLog)) { + amd::Comgr::destroy_action_info(action); + amd::Comgr::destroy_data_set(output); + return false; + } + + if (!extractByteCodeBinary(output, AMD_COMGR_DATA_KIND_BC, LLVMBitcode)) { + amd::Comgr::destroy_action_info(action); + amd::Comgr::destroy_data_set(output); + return false; + } + + // Clean up + amd::Comgr::destroy_action_info(action); + amd::Comgr::destroy_data_set(output); + return true; +} + +bool linkLLVMBitcode(const amd_comgr_data_set_t linkInputs, const std::string& isa, + std::vector& linkOptions, std::string& buildLog, + std::vector& LinkedLLVMBitcode) { + amd_comgr_language_t lang = AMD_COMGR_LANGUAGE_HIP; + amd_comgr_action_info_t action; + amd_comgr_data_set_t dataSetDevLibs; + + if (auto res = createAction(action, linkOptions, isa, AMD_COMGR_LANGUAGE_HIP); + res != AMD_COMGR_STATUS_SUCCESS) { + return false; + } + + if (auto res = amd::Comgr::create_data_set(&dataSetDevLibs); res != AMD_COMGR_STATUS_SUCCESS) { + amd::Comgr::destroy_action_info(action); + return false; + } + + + if (auto res = amd::Comgr::do_action(AMD_COMGR_ACTION_ADD_DEVICE_LIBRARIES, action, linkInputs, + dataSetDevLibs); + res != AMD_COMGR_STATUS_SUCCESS) { + extractBuildLog(dataSetDevLibs, buildLog); + LogPrintfInfo("%s", buildLog.c_str()); + amd::Comgr::destroy_action_info(action); + amd::Comgr::destroy_data_set(dataSetDevLibs); + return false; + } + + if (!extractBuildLog(dataSetDevLibs, buildLog)) { + amd::Comgr::destroy_action_info(action); + amd::Comgr::destroy_data_set(dataSetDevLibs); + return false; + } + + amd_comgr_data_set_t output; + if (auto res = amd::Comgr::create_data_set(&output); res != AMD_COMGR_STATUS_SUCCESS) { + amd::Comgr::destroy_action_info(action); + amd::Comgr::destroy_data_set(dataSetDevLibs); + return false; + } + + if (auto res = + amd::Comgr::do_action(AMD_COMGR_ACTION_LINK_BC_TO_BC, action, dataSetDevLibs, output); + res != AMD_COMGR_STATUS_SUCCESS) { + amd::Comgr::destroy_action_info(action); + amd::Comgr::destroy_data_set(dataSetDevLibs); + amd::Comgr::destroy_data_set(output); + return false; + } + + if (!extractBuildLog(output, buildLog)) { + amd::Comgr::destroy_action_info(action); + amd::Comgr::destroy_data_set(dataSetDevLibs); + amd::Comgr::destroy_data_set(output); + return false; + } + + if (!extractByteCodeBinary(output, AMD_COMGR_DATA_KIND_BC, LinkedLLVMBitcode)) { + amd::Comgr::destroy_action_info(action); + amd::Comgr::destroy_data_set(dataSetDevLibs); + amd::Comgr::destroy_data_set(output); + return false; + } + + amd::Comgr::destroy_action_info(action); + amd::Comgr::destroy_data_set(dataSetDevLibs); + amd::Comgr::destroy_data_set(output); + return true; +} + +bool createExecutable(const amd_comgr_data_set_t linkInputs, const std::string& isa, + std::vector& exeOptions, std::string& buildLog, + std::vector& executable) { + amd_comgr_action_info_t action; + + if (auto res = createAction(action, exeOptions, isa); res != AMD_COMGR_STATUS_SUCCESS) { + return false; + } + + amd_comgr_data_set_t relocatableData; + if (auto res = amd::Comgr::create_data_set(&relocatableData); res != AMD_COMGR_STATUS_SUCCESS) { + amd::Comgr::destroy_action_info(action); + return false; + } + + if (auto res = amd::Comgr::do_action(AMD_COMGR_ACTION_CODEGEN_BC_TO_RELOCATABLE, action, + linkInputs, relocatableData); + res != AMD_COMGR_STATUS_SUCCESS) { + extractBuildLog(relocatableData, buildLog); + amd::Comgr::destroy_action_info(action); + amd::Comgr::destroy_data_set(relocatableData); + return false; + } + + if (!extractBuildLog(relocatableData, buildLog)) { + amd::Comgr::destroy_action_info(action); + amd::Comgr::destroy_data_set(relocatableData); + return false; + } + + + amd::Comgr::destroy_action_info(action); + std::vector emptyOpt; + if (auto res = createAction(action, emptyOpt, isa); res != AMD_COMGR_STATUS_SUCCESS) { + amd::Comgr::destroy_data_set(relocatableData); + return false; + } + + amd_comgr_data_set_t output; + if (auto res = amd::Comgr::create_data_set(&output); res != AMD_COMGR_STATUS_SUCCESS) { + amd::Comgr::destroy_action_info(action); + amd::Comgr::destroy_data_set(relocatableData); + return false; + } + + if (auto res = amd::Comgr::do_action(AMD_COMGR_ACTION_LINK_RELOCATABLE_TO_EXECUTABLE, action, + relocatableData, output); + res != AMD_COMGR_STATUS_SUCCESS) { + extractBuildLog(output, buildLog); + amd::Comgr::destroy_action_info(action); + amd::Comgr::destroy_data_set(output); + amd::Comgr::destroy_data_set(relocatableData); + return false; + } + + if (!extractBuildLog(output, buildLog)) { + amd::Comgr::destroy_action_info(action); + amd::Comgr::destroy_data_set(output); + amd::Comgr::destroy_data_set(relocatableData); + return false; + } + + if (!extractByteCodeBinary(output, AMD_COMGR_DATA_KIND_EXECUTABLE, executable)) { + amd::Comgr::destroy_action_info(action); + amd::Comgr::destroy_data_set(output); + amd::Comgr::destroy_data_set(relocatableData); + return false; + } + + amd::Comgr::destroy_action_info(action); + amd::Comgr::destroy_data_set(output); + amd::Comgr::destroy_data_set(relocatableData); + + return true; +} + +bool dumpIsaFromBC(const amd_comgr_data_set_t isaInputs, const std::string& isa, + std::vector& exeOptions, std::string name, std::string& buildLog) { + if (name.size() == 0) return false; + amd_comgr_action_info_t action; + + if (auto res = createAction(action, exeOptions, isa); res != AMD_COMGR_STATUS_SUCCESS) { + return false; + } + + amd_comgr_data_set_t isaData; + if (auto res = amd::Comgr::create_data_set(&isaData); res != AMD_COMGR_STATUS_SUCCESS) { + amd::Comgr::destroy_action_info(action); + return false; + } + + if (auto res = amd::Comgr::do_action(AMD_COMGR_ACTION_CODEGEN_BC_TO_ASSEMBLY, action, isaInputs, + isaData); + res != AMD_COMGR_STATUS_SUCCESS) { + extractBuildLog(isaData, buildLog); + amd::Comgr::destroy_action_info(action); + amd::Comgr::destroy_data_set(isaData); + return false; + } + + std::vector isaOutput; + if (!extractByteCodeBinary(isaData, AMD_COMGR_DATA_KIND_SOURCE, isaOutput)) { + amd::Comgr::destroy_action_info(action); + amd::Comgr::destroy_data_set(isaData); + return false; + } + + auto isaFileName = name + ".s"; + std::ofstream f(isaFileName.c_str(), std::ios::trunc | std::ios::binary); + if (f.is_open()) { + f.write(isaOutput.data(), isaOutput.size()); + f.close(); + } else { + buildLog += "Warning: writing isa file failed.\n"; + amd::Comgr::destroy_action_info(action); + amd::Comgr::destroy_data_set(isaData); + return false; + } + amd::Comgr::destroy_action_info(action); + amd::Comgr::destroy_data_set(isaData); + return true; +} + +bool demangleName(const std::string& mangledName, std::string& demangledName) { + amd_comgr_data_t mangled_data; + amd_comgr_data_t demangled_data; + + if (AMD_COMGR_STATUS_SUCCESS != amd::Comgr::create_data(AMD_COMGR_DATA_KIND_BYTES, &mangled_data)) + return false; + + if (AMD_COMGR_STATUS_SUCCESS != + amd::Comgr::set_data(mangled_data, mangledName.size(), mangledName.c_str())) { + amd::Comgr::release_data(mangled_data); + return false; + } + + if (AMD_COMGR_STATUS_SUCCESS != amd::Comgr::demangle_symbol_name(mangled_data, &demangled_data)) { + amd::Comgr::release_data(mangled_data); + return false; + } + + size_t demangled_size = 0; + if (AMD_COMGR_STATUS_SUCCESS != amd::Comgr::get_data(demangled_data, &demangled_size, NULL)) { + amd::Comgr::release_data(mangled_data); + amd::Comgr::release_data(demangled_data); + return false; + } + + demangledName.resize(demangled_size); + + if (AMD_COMGR_STATUS_SUCCESS != + amd::Comgr::get_data(demangled_data, &demangled_size, + const_cast(demangledName.data()))) { + amd::Comgr::release_data(mangled_data); + amd::Comgr::release_data(demangled_data); + return false; + } + + amd::Comgr::release_data(mangled_data); + amd::Comgr::release_data(demangled_data); + return true; +} + +std::string handleMangledName(std::string loweredName) { + if (loweredName.empty()) { + return loweredName; + } + + if (loweredName.find(".kd") != std::string::npos) { + return {}; + } + + if (loweredName.find("void ") == 0) { + loweredName.erase(0, strlen("void ")); + } + + auto dx{loweredName.find_first_of("(<")}; + + if (dx == std::string::npos) { + return loweredName; + } + + if (loweredName[dx] == '<') { + uint32_t count = 1; + do { + ++dx; + count += (loweredName[dx] == '<') ? 1 : ((loweredName[dx] == '>') ? -1 : 0); + } while (count); + + loweredName.erase(++dx); + } else { + loweredName.erase(dx); + } + + return loweredName; +} + +bool fillDemangledNames(std::vector& executable, std::vector& mangledNames) { + amd_comgr_data_t dataObject; + if (auto res = amd::Comgr::create_data(AMD_COMGR_DATA_KIND_EXECUTABLE, &dataObject); + res != AMD_COMGR_STATUS_SUCCESS) { + return false; + } + + if (auto res = amd::Comgr::set_data(dataObject, executable.size(), executable.data())) { + amd::Comgr::release_data(dataObject); + return false; + } + + auto callback = [](amd_comgr_symbol_t symbol, void* data) { + if (data == nullptr) return AMD_COMGR_STATUS_ERROR_INVALID_ARGUMENT; + size_t len = 0; + if (auto res = amd::Comgr::symbol_get_info(symbol, AMD_COMGR_SYMBOL_INFO_NAME_LENGTH, &len); + res != AMD_COMGR_STATUS_SUCCESS) + return res; + std::string name(len, 0); + if (auto res = amd::Comgr::symbol_get_info(symbol, AMD_COMGR_SYMBOL_INFO_NAME, &name[0]); + res != AMD_COMGR_STATUS_SUCCESS) + return res; + auto storage = reinterpret_cast*>(data); + storage->push_back(name); + return AMD_COMGR_STATUS_SUCCESS; + }; + + if (auto res = + amd::Comgr::iterate_symbols(dataObject, callback, reinterpret_cast(&mangledNames)); + res != AMD_COMGR_STATUS_SUCCESS) { + amd::Comgr::release_data(dataObject); + return false; + } + + amd::Comgr::release_data(dataObject); + return true; +} + +bool getMangledNames(const std::vector& mangledNames, + std::map& strippedNames, + std::map& demangledNames) { + for (auto& i : mangledNames) { + std::string demangledName; + if (!demangleName(i, demangledName)) return false; + demangledName = handleMangledName(demangledName); + + demangledName.erase(std::remove_if(demangledName.begin(), demangledName.end(), + [](unsigned char c) { return std::isspace(c); }), + demangledName.end()); + + if (auto res = strippedNames.find(demangledName); res != strippedNames.end()) { + auto& strippedName = res->second; + if (auto dres = demangledNames.find(strippedName); dres != demangledNames.end()) { + dres->second = i; + continue; + } else { + return false; + } + } + + if (auto dres = demangledNames.find(demangledName); dres != demangledNames.end()) { + dres->second = i; + continue; + } + } + + return true; +} +} // namespace helpers +} // namespace hiprtc diff --git a/projects/clr/hipamd/src/hiprtc/hiprtcComgrHelper.hpp b/projects/clr/hipamd/src/hiprtc/hiprtcComgrHelper.hpp new file mode 100644 index 0000000000..f9d403069a --- /dev/null +++ b/projects/clr/hipamd/src/hiprtc/hiprtcComgrHelper.hpp @@ -0,0 +1,60 @@ +/* +Copyright (c) 2022 - Present 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. +*/ + +#pragma once + +#include +#include + +#include "vdi_common.hpp" +#include "utils/debug.hpp" +#include "device/comgrctx.hpp" + +namespace hiprtc { +namespace helpers { +bool addCodeObjData(amd_comgr_data_set_t& input, const std::vector& source, + const std::string& name, const amd_comgr_data_kind_t type); +bool extractBuildLog(amd_comgr_data_set_t dataSet, std::string& buildLog); +bool extractByteCodeBinary(const amd_comgr_data_set_t inDataSet, + const amd_comgr_data_kind_t dataKind, std::vector& bin); +bool createAction(amd_comgr_action_info_t& action, std::vector& options, + const std::string& isa, + const amd_comgr_language_t lang = AMD_COMGR_LANGUAGE_NONE); +bool compileToBitCode(const amd_comgr_data_set_t compileInputs, const std::string& isa, + std::vector& compileOptions, std::string& buildLog, + std::vector& LLVMBitcode); +bool linkLLVMBitcode(const amd_comgr_data_set_t linkInputs, const std::string& isa, + std::vector& linkOptions, std::string& buildLog, + std::vector& LinkedLLVMBitcode); +bool createExecutable(const amd_comgr_data_set_t linkInputs, const std::string& isa, + std::vector& exeOptions, std::string& buildLog, + std::vector& executable); +bool dumpIsaFromBC(const amd_comgr_data_set_t isaInputs, const std::string& isa, + std::vector& exeOptions, std::string name, std::string& buildLog); +bool demangleName(const std::string& mangledName, std::string& demangledName); +std::string handleMangledName(std::string loweredName); +bool fillDemangledNames(std::vector& executable, std::vector& mangledNames); +bool getMangledNames(const std::vector& mangledNames, + std::map& strippedNames, + std::map& demangledNames); +} // namespace helpers +} // namespace hiprtc diff --git a/projects/clr/hipamd/src/hiprtc/hiprtcInternal.cpp b/projects/clr/hipamd/src/hiprtc/hiprtcInternal.cpp new file mode 100644 index 0000000000..d0783d092f --- /dev/null +++ b/projects/clr/hipamd/src/hiprtc/hiprtcInternal.cpp @@ -0,0 +1,295 @@ +/* +Copyright (c) 2022 - Present 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. +*/ + +#include "hiprtcInternal.hpp" + +#include "vdi_common.hpp" +#include "utils/flags.hpp" + +namespace hiprtc { +using namespace helpers; +RTCProgram::RTCProgram(std::string name_) : name(name_) { + std::call_once(amd::Comgr::initialized, amd::Comgr::LoadLib); + if (amd::Comgr::create_data_set(&compileInput) != AMD_COMGR_STATUS_SUCCESS || + amd::Comgr::create_data_set(&linkInput) != AMD_COMGR_STATUS_SUCCESS || + amd::Comgr::create_data_set(&execInput) != AMD_COMGR_STATUS_SUCCESS) { + crashWithMessage("Failed to allocate internal hiprtc structure"); + } + + // Add internal header + if (!addBuiltinHeader()) { + crashWithMessage("Unable to add internal header"); + } + + // Add compile options + const std::string hipVerOpt{"--hip-version=" + std::to_string(HIP_VERSION_MAJOR) + '.' + + std::to_string(HIP_VERSION_MINOR) + '.' + + std::to_string(HIP_VERSION_PATCH)}; + const std::string hipVerMajor{"-DHIP_VERSION_MAJOR=" + std::to_string(HIP_VERSION_MAJOR)}; + const std::string hipVerMinor{"-DHIP_VERSION_MINOR=" + std::to_string(HIP_VERSION_MINOR)}; + const std::string hipVerPatch{"-DHIP_VERSION_PATCH=" + std::to_string(HIP_VERSION_PATCH)}; + + compileOptions.reserve(18); // count of options below + compileOptions.push_back("-O3"); + +#ifdef HIPRTC_EARLY_INLINE + compileOptions.push_back("-mllvm"); + compileOptions.push_back("-amdgpu-early-inline-all"); +#endif + compileOptions.push_back("-mllvm"); + compileOptions.push_back("-amdgpu-prelink"); + + if (GPU_ENABLE_WGP_MODE) compileOptions.push_back("-mcumode"); + + if (!GPU_ENABLE_WAVE32_MODE) compileOptions.push_back("-mwavefrontsize64"); + + compileOptions.push_back(hipVerOpt); + compileOptions.push_back(hipVerMajor); + compileOptions.push_back(hipVerMinor); + compileOptions.push_back(hipVerPatch); + compileOptions.push_back("-D__HIPCC_RTC__"); + compileOptions.push_back("-include"); + compileOptions.push_back("hiprtc_runtime.h"); + compileOptions.push_back("-std=c++14"); + compileOptions.push_back("-nogpuinc"); +#ifdef _WIN32 + compileOptions.push_back("-target x86_64-pc-windows-msvc"); + compileOptions.push_back("-fms-extensions"); + compileOptions.push_back("-fms-compatibility"); +#endif + + if (!GPU_ENABLE_WAVE32_MODE) linkOptions.push_back("wavefrontsize64"); + + exeOptions.push_back("-O3"); + exeOptions.push_back("-mllvm"); + exeOptions.push_back("-amdgpu-internalize-symbols"); + exeOptions.push_back("-mcumode"); + if (!GPU_ENABLE_WAVE32_MODE) exeOptions.push_back("-mwavefrontsize64"); +} + +bool RTCProgram::addSource(const std::string& source, const std::string& name) { + if (source.size() == 0 || name.size() == 0) { + LogError("Error in hiprtc: source or name is of size 0 in addSource"); + return false; + } + sourceCode += source; + sourceName = name; + return true; +} + +// addSource_impl is a different function because we need to add source when we track mangled +// objects +bool RTCProgram::addSource_impl() { + std::vector vsource(sourceCode.begin(), sourceCode.end()); + if (!addCodeObjData(compileInput, vsource, sourceName, AMD_COMGR_DATA_KIND_SOURCE)) { + return false; + } + return true; +} + +bool RTCProgram::addHeader(const std::string& source, const std::string& name) { + if (source.size() == 0 || name.size() == 0) { + LogError("Error in hiprtc: source or name is of size 0 in addHeader"); + return false; + } + std::vector vsource(source.begin(), source.end()); + if (!addCodeObjData(compileInput, vsource, name, AMD_COMGR_DATA_KIND_INCLUDE)) { + return false; + } + return true; +} + +bool RTCProgram::addBuiltinHeader() { + std::vector source(__hipRTC_header, __hipRTC_header + __hipRTC_header_size); + std::string name{"hiprtc_runtime.h"}; + if (!addCodeObjData(compileInput, source, name, AMD_COMGR_DATA_KIND_INCLUDE)) { + return false; + } + return true; +} + +bool RTCProgram::transformOptions() { + auto getValueOf = [](const std::string& option) { + std::string res; + auto f = std::find(option.begin(), option.end(), '='); + if (f != option.end()) res = std::string(f + 1, option.end()); + return res; + }; + + for (auto& i : compileOptions) { + if (i == "-hip-pch") { + LogInfo( + "-hip-pch is deprecated option, has no impact on execution of new hiprtc programs, it " + "can be removed"); + i.clear(); + continue; + } + // Some rtc samples use --gpu-architecture + if (i.rfind("--gpu-architecture=", 0) == 0) { + LogInfo("--gpu-architecture is nvcc option, transforming it to --offload-arch option"); + auto val = getValueOf(i); + i = "--offload-arch=" + val; + continue; + } + if (i == "--save-temps") { + settings.dumpISA = true; + continue; + } + } + + if (auto res = std::find_if( + compileOptions.begin(), compileOptions.end(), + [](const std::string& str) { return str.find("--offload-arch=") != std::string::npos; }); + res != compileOptions.end()) { + auto isaName = getValueOf(*res); + isa = "amdgcn-amd-amdhsa--" + isaName; + settings.offloadArchProvided = true; + return true; + } + buildLog += + "Error: Please provide architecture for which code is to be " + "generated.\n"; + return false; +} + +amd::Monitor RTCProgram::lock_("HIPRTC Program", true); + +bool RTCProgram::compile(const std::vector& options) { + amd::ScopedLock lock(lock_); // Lock, because LLVM is not multi threaded + + if (!addSource_impl()) { + LogError("Error in hiprtc: unable to add source code"); + return false; + } + + // Append compile options + compileOptions.reserve(compileOptions.size() + options.size()); + compileOptions.insert(compileOptions.end(), options.begin(), options.end()); + + if (!transformOptions()) { + LogError("Error in hiprtc: unable to transform options"); + return false; + } + + std::vector LLVMBitcode; + if (!compileToBitCode(compileInput, isa, compileOptions, buildLog, LLVMBitcode)) { + LogError("Error in hiprtc: unable to compile source to bitcode"); + return false; + } + + std::string linkFileName = "linked"; + if (!addCodeObjData(linkInput, LLVMBitcode, linkFileName, AMD_COMGR_DATA_KIND_BC)) { + LogError("Error in hiprtc: unable to add linked code object"); + return false; + } + + std::vector LinkedLLVMBitcode; + if (!linkLLVMBitcode(linkInput, isa, linkOptions, buildLog, LinkedLLVMBitcode)) { + LogError("Error in hiprtc: unable to add device libs to linked bitcode"); + return false; + } + + std::string linkedFileName = "LLVMBitcode.bc"; + if (!addCodeObjData(execInput, LinkedLLVMBitcode, linkedFileName, AMD_COMGR_DATA_KIND_BC)) { + LogError("Error in hiprtc: unable to add device libs linked code object"); + return false; + } + + if (settings.dumpISA) { + if (!dumpIsaFromBC(execInput, isa, exeOptions, name, buildLog)) { + LogError("Error in hiprtc: unable to dump isa code"); + return false; + } + } + + if (!createExecutable(execInput, isa, exeOptions, buildLog, executable)) { + LogError("Error in hiprtc: unable to create executable"); + return false; + } + + std::vector mangledNames; + if (!fillDemangledNames(executable, mangledNames)) { + LogError("Error in hiprtc: unable to fill demangled names"); + return false; + } + + if (!getMangledNames(mangledNames, strippedNames, demangledNames)) { + LogError("Error in hiprtc: unable to get mangled names"); + return false; + } + + return true; +} + +bool RTCProgram::trackMangledName(std::string& name) { + amd::ScopedLock lock(lock_); + + if (name.size() == 0) return false; + + std::string strippedName = name; + if (strippedName.back() == ')') { + strippedName.pop_back(); + strippedName.erase(0, strippedName.find('(')); + } + if (strippedName.front() == '&') { + strippedName.erase(0, 1); + } + + std::string strippedNameNoSpace = strippedName; + strippedNameNoSpace.erase(std::remove_if(strippedNameNoSpace.begin(), + strippedNameNoSpace.end(), + [](unsigned char c) { + return std::isspace(c); + }), strippedNameNoSpace.end()); + + strippedNames.insert(std::pair(name, strippedNameNoSpace)); + demangledNames.insert(std::pair(strippedName, "")); + + const auto var{"__hiprtc_" + std::to_string(strippedNames.size())}; + const auto code{"\nextern \"C\" constexpr auto " + var + " = " + name + ";\n"}; + + sourceCode += code; + return true; +} + +bool RTCProgram::getDemangledName(const char* name_expression, const char** loweredName) { + std::string name = name_expression; + if (auto res = strippedNames.find(name); res != strippedNames.end()) { + if (auto dres = demangledNames.find(res->second); dres != demangledNames.end()) { + if (dres->second.size() != 0) { + *loweredName = dres->second.c_str(); + return true; + } else + return false; + } + } + if (auto dres = demangledNames.find(name); dres != demangledNames.end()) { + if (dres->second.size() != 0) { + *loweredName = dres->second.c_str(); + return true; + } + return false; + } + return false; +} + +} // namespace hiprtc diff --git a/projects/clr/hipamd/src/hiprtc/hiprtcInternal.hpp b/projects/clr/hipamd/src/hiprtc/hiprtcInternal.hpp new file mode 100644 index 0000000000..9667ebb8eb --- /dev/null +++ b/projects/clr/hipamd/src/hiprtc/hiprtcInternal.hpp @@ -0,0 +1,155 @@ +/* +Copyright (c) 2022 - Present 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. +*/ + +#include +#include + + +#ifdef HIPRTC_USE_EXCEPTIONS +#include +#endif +#include +#include +#include +#include + +#include "top.hpp" +#include "utils/debug.hpp" +#include "utils/flags.hpp" +#include "utils/macros.hpp" + +#ifdef __HIP_ENABLE_RTC +extern "C" { +extern const char __hipRTC_header[]; +extern unsigned __hipRTC_header_size; +} +#endif + +#include "hiprtcComgrHelper.hpp" + + +namespace hiprtc { +namespace internal { +template inline std::string ToString(T v) { + std::ostringstream ss; + ss << v; + return ss.str(); +} + +inline std::string ToString() { return (""); } + +template inline std::string ToString(T first, Args... args) { + return ToString(first) + ", " + ToString(args...); +} +} // namespace internal +} // namespace hiprtc + +#define HIPRTC_INIT_API(...) \ + ClPrint(amd::LOG_INFO, amd::LOG_API, "%s ( %s )", __func__, \ + hiprtc::internal::ToString(__VA_ARGS__).c_str()); + +#define HIPRTC_RETURN(ret) \ + hiprtc::g_lastRtcError = (ret); \ + ClPrint(amd::LOG_INFO, amd::LOG_API, "%s: Returned %s", __func__, \ + hiprtcGetErrorString(hiprtc::g_lastRtcError)); \ + return hiprtc::g_lastRtcError; + + +namespace hiprtc { + +static void crashWithMessage(std::string message) { +#ifdef HIPRTC_USE_EXCEPTIONS + throw std::runtime_error(message); +#else + guarantee(false, message.c_str()); +#endif +} + +struct Settings { + bool dumpISA{false}; + bool offloadArchProvided{false}; +}; + +class RTCProgram { + static amd::Monitor lock_; + static std::once_flag initialized; + + std::string name; + Settings settings; + + std::string isa; + std::string buildLog; + + std::vector executable; + + std::map strippedNames; + std::map demangledNames; + std::string sourceCode; + std::string sourceName; + + std::vector compileOptions; + std::vector linkOptions; + std::vector exeOptions; + + amd_comgr_data_set_t compileInput; + amd_comgr_data_set_t linkInput; + amd_comgr_data_set_t execInput; + + bool dumpIsa(); + + bool addSource_impl(); + bool addBuiltinHeader(); + bool transformOptions(); + + RTCProgram() = delete; + RTCProgram(RTCProgram&) = delete; + RTCProgram& operator=(RTCProgram&) = delete; + + public: + RTCProgram(std::string); + + // Converters + inline static hiprtcProgram as_hiprtcProgram(RTCProgram* p) { + return reinterpret_cast(p); + } + inline static RTCProgram* as_RTCProgram(hiprtcProgram& p) { + return reinterpret_cast(p); + } + + bool addSource(const std::string& source, const std::string& name); + bool addHeader(const std::string& source, const std::string& name); + bool compile(const std::vector& options); + bool getDemangledName(const char* name_expression, const char** loweredName); + bool trackMangledName(std::string& name); + + const std::vector& getExec() const { return executable; } + size_t getExecSize() const { return executable.size(); } + const std::string& getLog() const { return buildLog; } + size_t getLogSize() const { return buildLog.size(); } + + ~RTCProgram() { + amd::Comgr::destroy_data_set(compileInput); + amd::Comgr::destroy_data_set(linkInput); + amd::Comgr::destroy_data_set(execInput); + } +}; +} // namespace hiprtc