AI/rocm-systems
2
0
Форкнуть 0
Код Задачи Запросы на слияние Действия Пакеты Проекты Релизы Вики Активность

SWDEV-308763 - Make hiprtc a separate lib and rewrite it using comgr

Просмотр исходного кода 
Change-Id: I28ac50ec897accb2a2b4590d3c25965d907426fb


[ROCm/clr commit: e6c0086d1c]
Этот коммит содержится в:
cjatin
2021-10-01 01:46:34 +05:30
коммит произвёл Satyanvesh Dittakavi
родитель 8139cc4683
Коммит added78cbc
14 изменённых файлов: 1518 добавлений и 95 удалений
Скачать Patch файл Скачать Diff файл Показать все файлы Свернуть все файлы
projects/clr/hipamd/CMakeLists.txt
-1
Просмотреть файл
@@ -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")
projects/clr/hipamd/packaging/CMakeLists.txt
+6 -5
Просмотреть файл
@@ -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
projects/clr/hipamd/src/CMakeLists.txt
+13 -78
Просмотреть файл
@@ -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 $<TARGET_FILE:clang> -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 $<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})
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 $<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)
install(TARGETS hiprtc-builtins
RUNTIME DESTINATION ${CMAKE_INSTALL_BINDIR}
LIBRARY DESTINATION lib
ARCHIVE DESTINATION ${CMAKE_INSTALL_LIBDIR})
endif()
#############################
projects/clr/hipamd/src/amdhip.def
-10
Просмотреть файл
@@ -239,16 +239,6 @@ hipDestroySurfaceObject
hipInitActivityCallback
hipEnableActivityCallback
hipGetCmdName
hiprtcAddNameExpression
hiprtcCompileProgram
hiprtcCreateProgram
hiprtcDestroyProgram
hiprtcGetLoweredName
hiprtcGetProgramLog
hiprtcGetProgramLogSize
hiprtcGetCode
hiprtcGetCodeSize
hiprtcGetErrorString
hipMipmappedArrayCreate
hipMallocMipmappedArray
hipMipmappedArrayDestroy
projects/clr/hipamd/src/hiprtc/CMakeLists.txt
+195
Просмотреть файл
@@ -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} $<TARGET_FILE:hiprtc>)
endif()
else()
add_library(hiprtc STATIC $<TARGET_OBJECTS:rocclr>)
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_PROPERTY:rocclr,COMPILE_DEFINITIONS>)
target_include_directories(hiprtcobject PRIVATE $<TARGET_PROPERTY:rocclr,INCLUDE_DIRECTORIES>)
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 $<TARGET_FILE:clang> -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 $<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})
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})
projects/clr/hipamd/src/cmake/HIPRTC.cmake → projects/clr/hipamd/src/hiprtc/cmake/HIPRTC.cmake
+3 -1
Просмотреть файл
@@ -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\
projects/clr/hipamd/src/hiprtc-builtins.def → projects/clr/hipamd/src/hiprtc/hiprtc-builtins.def
Просмотреть файл
projects/clr/hipamd/src/hiprtc/hiprtc.cpp
+226
Просмотреть файл
@@ -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 <hip/hiprtc.h>
#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<std::string> 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);
}
projects/clr/hipamd/src/hiprtc/hiprtc.def
+11
Просмотреть файл
@@ -0,0 +1,11 @@
EXPORTS
hiprtcAddNameExpression
hiprtcCompileProgram
hiprtcCreateProgram
hiprtcDestroyProgram
hiprtcGetLoweredName
hiprtcGetProgramLog
hiprtcGetProgramLogSize
hiprtcGetCode
hiprtcGetCodeSize
hiprtcGetErrorString
projects/clr/hipamd/src/hiprtc/hiprtc.map.in
+16
Просмотреть файл
@@ -0,0 +1,16 @@
{
global:
hiprtcCompileProgram;
hiprtcCreateProgram;
hiprtcDestroyProgram;
hiprtcGetLoweredName;
hiprtcGetProgramLog;
hiprtcGetProgramLogSize;
hiprtcGetCode;
hiprtcGetCodeSize;
hiprtcGetErrorString;
hiprtcAddNameExpression;
hiprtcVersion;
local:
*;
};
projects/clr/hipamd/src/hiprtc/hiprtcComgrHelper.cpp
+538
Просмотреть файл
@@ -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<char>& 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<char> 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<char>& 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<std::string>& 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<const char*> 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<std::string>& compileOptions, std::string& buildLog,
std::vector<char>& 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<std::string>& linkOptions, std::string& buildLog,
std::vector<char>& 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<std::string>& exeOptions, std::string& buildLog,
std::vector<char>& 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<std::string> 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<std::string>& 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<char> 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<char*>(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<char>& executable, std::vector<std::string>& 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<std::vector<std::string>*>(data);
storage->push_back(name);
return AMD_COMGR_STATUS_SUCCESS;
};
if (auto res =
amd::Comgr::iterate_symbols(dataObject, callback, reinterpret_cast<void*>(&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<std::string>& mangledNames,
std::map<std::string, std::string>& strippedNames,
std::map<std::string, std::string>& 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
projects/clr/hipamd/src/hiprtc/hiprtcComgrHelper.hpp
+60
Просмотреть файл
@@ -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 <vector>
#include <string>
#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<char>& 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<char>& bin);
bool createAction(amd_comgr_action_info_t& action, std::vector<std::string>& 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<std::string>& compileOptions, std::string& buildLog,
std::vector<char>& LLVMBitcode);
bool linkLLVMBitcode(const amd_comgr_data_set_t linkInputs, const std::string& isa,
std::vector<std::string>& linkOptions, std::string& buildLog,
std::vector<char>& LinkedLLVMBitcode);
bool createExecutable(const amd_comgr_data_set_t linkInputs, const std::string& isa,
std::vector<std::string>& exeOptions, std::string& buildLog,
std::vector<char>& executable);
bool dumpIsaFromBC(const amd_comgr_data_set_t isaInputs, const std::string& isa,
std::vector<std::string>& 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<char>& executable, std::vector<std::string>& mangledNames);
bool getMangledNames(const std::vector<std::string>& mangledNames,
std::map<std::string, std::string>& strippedNames,
std::map<std::string, std::string>& demangledNames);
} // namespace helpers
} // namespace hiprtc
projects/clr/hipamd/src/hiprtc/hiprtcInternal.cpp
+295
Просмотреть файл
@@ -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<char> 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<char> vsource(source.begin(), source.end());
if (!addCodeObjData(compileInput, vsource, name, AMD_COMGR_DATA_KIND_INCLUDE)) {
return false;
}
return true;
}
bool RTCProgram::addBuiltinHeader() {
std::vector<char> 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<std::string>& 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<char> 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<char> 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<std::string> 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<std::string, std::string>(name, strippedNameNoSpace));
demangledNames.insert(std::pair<std::string, std::string>(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
projects/clr/hipamd/src/hiprtc/hiprtcInternal.hpp
+155
Просмотреть файл
@@ -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 <hip/hiprtc.h>
#include <hip/hip_version.h>
#ifdef HIPRTC_USE_EXCEPTIONS
#include <exception>
#endif
#include <atomic>
#include <map>
#include <mutex>
#include <string>
#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 <typename T> inline std::string ToString(T v) {
std::ostringstream ss;
ss << v;
return ss.str();
}
inline std::string ToString() { return (""); }
template <typename T, typename... Args> 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<char> executable;
std::map<std::string, std::string> strippedNames;
std::map<std::string, std::string> demangledNames;
std::string sourceCode;
std::string sourceName;
std::vector<std::string> compileOptions;
std::vector<std::string> linkOptions;
std::vector<std::string> 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<hiprtcProgram>(p);
}
inline static RTCProgram* as_RTCProgram(hiprtcProgram& p) {
return reinterpret_cast<RTCProgram*>(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<std::string>& options);
bool getDemangledName(const char* name_expression, const char** loweredName);
bool trackMangledName(std::string& name);
const std::vector<char>& 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