AI/rocm-systems
2
0
Forka 0
Codice Problemi Pull Request Actions Pacchetti Progetti Rilasci Wiki Attività

Update to modern cmake.

Sfoglia il codice sorgente 
Adds support for find_package(), locates dependencies with
find_package(), swaps the roles of /hsa/include/hsa and /include/hsa
as well as /lib & /hsa/lib.

Kernel code objects no longer build at every make call but only
as needed.  Dependencies are tracked through to clang.

Device lib is still located with directory searches.  build_devicelibs.sh
does not yet install the cmake config files on the build systems.

Corrects DAZ mode mismatch in code object compilation.

Still needs updating to compiler properties rather than direct
manipulation of CMAKE_CXX_FLAGS.

Change-Id: I02d946c8a77d5cf753681f8e3d3153fca4aae86a


[ROCm/ROCR-Runtime commit: 55a4f01b16]
This commit is contained in:
Sean Keely
2020-05-26 19:17:14 -05:00
parent 960fcd7a41
commit 40ac4daa21
10 ha cambiato i file con 315 aggiunte e 872 eliminazioni
Scarica il file Patch Scarica il file Diff Expand all files Collapse all files
projects/rocr-runtime/runtime/hsa-runtime/CMakeLists.txt
+204 -114
Vedi File
@@ -47,16 +47,18 @@ set ( CORE_RUNTIME_NAME "hsa-runtime" )
set ( CORE_RUNTIME_TARGET "${CORE_RUNTIME_NAME}64" )
set ( CORE_RUNTIME_LIBRARY "lib${CORE_RUNTIME_TARGET}" )
## Set image module name
set ( IMAGE_NAME "hsa-ext-image" )
set ( IMAGE_TARGET "${IMAGE_NAME}64" )
set ( IMAGE_LIBRARY "lib${IMAGE_TARGET}" )
## Set project name
project( ${CORE_RUNTIME_TARGET} )
## Utilty functions
list ( APPEND CMAKE_MODULE_PATH "${CMAKE_CURRENT_SOURCE_DIR}/cmake_modules" )
include ( utils )
include ( hsa_common )
include ( GNUInstallDirs )
## Expose static library option
if ( NOT DEFINED BUILD_SHARED_LIBS )
set ( BUILD_SHARED_LIBS "on" )
set ( BUILD_SHARED_LIBS ON )
endif()
set ( BUILD_SHARED_LIBS ${BUILD_SHARED_LIBS} CACHE BOOL "Build shared library (.so) or not.")
@@ -71,160 +73,249 @@ if (ROCM_CCACHE_BUILD)
endif() # if (CCACHE_PROGRAM)
endif() # if (ROCM_CCACHE_BUILD)
list ( APPEND CMAKE_MODULE_PATH "${CMAKE_CURRENT_SOURCE_DIR}/cmake_modules" )
include ( utils )
include ( hsa_common )
## Get version strings
get_version ( "1.2.0" )
if ( ${ROCM_PATCH_VERSION} )
set ( VERSION_PATCH ${ROCM_PATCH_VERSION})
endif()
set ( SO_VERSION_STRING "${VERSION_MAJOR}.${VERSION_MINOR}.${VERSION_PATCH}" )
set ( PACKAGE_VERSION_STRING "${VERSION_MAJOR}.${VERSION_MINOR}.${VERSION_PATCH}.${VERSION_COMMIT_COUNT}-${VERSION_JOB}-${VERSION_HASH}" )
## Find LibElf
## Find external dependencies.
find_package(LibElf REQUIRED)
find_package(hsakmt 1.0 REQUIRED)
## Create the rocr target.
add_library( ${CORE_RUNTIME_TARGET} "" )
## Check for memfd_create syscall
include(CheckSymbolExists)
CHECK_SYMBOL_EXISTS ( "__NR_memfd_create" "sys/syscall.h" HAVE_MEMFD_CREATE )
## Compiler preproc definitions.
add_definitions ( -D__linux__ )
add_definitions ( -DHSA_EXPORT=1 )
add_definitions ( -DHSA_EXPORT_FINALIZER=1 )
add_definitions ( -DHSA_EXPORT_IMAGES=1 )
add_definitions ( -D HSA_DEPRECATED= )
target_compile_definitions(${CORE_RUNTIME_TARGET} PRIVATE __linux__ HSA_EXPORT=1 HSA_EXPORT_FINALIZER=1 HSA_EXPORT_IMAGES=1 HSA_DEPRECATED=
ROCR_BUILD_ID=${PACKAGE_VERSION_STRING} )
## Image definitons - audit!
target_compile_definitions(${CORE_RUNTIME_TARGET} PRIVATE
UNIX_OS
LINUX
__AMD64__
__x86_64__
AMD_INTERNAL_BUILD
LITTLEENDIAN_CPU=1
BRAHMA_BUILD=1 )
if ( HAVE_MEMFD_CREATE )
add_definitions ( -DHAVE_MEMFD_CREATE )
target_compile_definitions(${CORE_RUNTIME_TARGET} PRIVATE HAVE_MEMFD_CREATE )
endif()
## Get the package version.
get_version ( "1.1.9" )
set (SO_MAJOR 1)
set (SO_MINOR 1)
if ( ${ROCM_PATCH_VERSION} )
set ( SO_PATCH ${ROCM_PATCH_VERSION})
set ( VERSION_PATCH ${ROCM_PATCH_VERSION})
else ()
set(SO_PATCH 9)
endif ()
set ( SO_VERSION_STRING "${SO_MAJOR}.${SO_MINOR}.${SO_PATCH}" )
set ( PACKAGE_VERSION_STRING "${VERSION_MAJOR}.${VERSION_MINOR}.${VERSION_PATCH}.${VERSION_COMMIT_COUNT}-${VERSION_JOB}-${VERSION_HASH}" )
## Find the hsakmt library and include files, use directory hint from cache
## Search relative to build directory, relative to source directory, and finally the rocm install default (/opt/rocm)
get_include_path( HSAKMT_INC_PATH "libhsakmt include path" NAMES "hsakmt.h" "libhsakmt/hsakmt.h" HINTS "${CMAKE_BINARY_DIR}/../../include" "${CMAKE_CURRENT_SOURCE_DIR}/../../../../libhsakmt/include" PATHS "/opt/rocm/include")
get_library_path( HSAKMT_LIB_PATH "libhsakmt library path" NAMES "libhsakmt.so" HINTS "${CMAKE_BINARY_DIR}/../../lib" "${CMAKE_BINARY_DIR}/../roct" PATHS "/opt/rocm/lib")
include_directories ( ${HSAKMT_INC_PATH} )
link_directories ( ${HSAKMT_LIB_PATH} )
## Set include directories for ROCr runtime
include_directories ( ${CMAKE_CURRENT_SOURCE_DIR} )
include_directories ( ${CMAKE_CURRENT_SOURCE_DIR}/libamdhsacode )
## ROCr build internal versioning
add_definitions ( -DROCR_BUILD_ID=${PACKAGE_VERSION_STRING} )
target_include_directories( ${CORE_RUNTIME_TARGET}
PUBLIC
$<BUILD_INTERFACE:${CMAKE_CURRENT_SOURCE_DIR}/inc>
$<INSTALL_INTERFACE:include/hsa>
PRIVATE
${CMAKE_CURRENT_SOURCE_DIR}
${CMAKE_CURRENT_SOURCE_DIR}/libamdhsacode )
## Set RUNPATH - ../../lib covers use of the legacy symlink in /hsa/lib/
set(CMAKE_INSTALL_RPATH "$ORIGIN;$ORIGIN/../../lib;$ORIGIN/../../lib64;$ORIGIN/../lib64")
set_property(TARGET ${CORE_RUNTIME_TARGET} PROPERTY INSTALL_RPATH "$ORIGIN;$ORIGIN/../../lib;$ORIGIN/../../lib64;$ORIGIN/../lib64" )
## ------------------------- Linux Compiler and Linker options -------------------------
set ( CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -Werror -fexceptions -fno-rtti -fvisibility=hidden -Wno-error=sign-compare -Wno-sign-compare -Wno-write-strings -Wno-conversion-null -fno-math-errno -fno-threadsafe-statics -fmerge-all-constants -fms-extensions -Wno-error=comment -Wno-comment -Wno-error=pointer-arith -Wno-pointer-arith -Wno-error=unused-variable -Wno-error=unused-function" )
## Extra image settings - audit!
set ( CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -Wno-deprecated-declarations -fPIC" )
if ( CMAKE_COMPILER_IS_GNUCXX )
set ( CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -Wno-error=unused-but-set-variable")
endif ()
set ( DRVDEF "${CMAKE_CURRENT_SOURCE_DIR}/hsacore.so.def" )
set ( LNKSCR "${CMAKE_CURRENT_SOURCE_DIR}/hsacore.so.link" )
set ( CMAKE_SHARED_LINKER_FLAGS "-Wl,-Bdynamic -Wl,-z,noexecstack -Wl,${LNKSCR} -Wl,--version-script=${DRVDEF} -Wl,--enable-new-dtags" )
## ------------------------- End Compiler and Linker options ----------------------------
## Source files.
set ( SRCS "core/util/lnx/os_linux.cpp"
"core/util/small_heap.cpp"
"core/util/timer.cpp"
"core/runtime/amd_blit_kernel.cpp"
"core/runtime/amd_blit_sdma.cpp"
"core/runtime/amd_cpu_agent.cpp"
"core/runtime/amd_gpu_agent.cpp"
"core/runtime/amd_hsa_loader.cpp"
"core/runtime/amd_aql_queue.cpp"
"core/runtime/amd_loader_context.cpp"
"core/runtime/hsa_ven_amd_loader.cpp"
"core/runtime/amd_memory_region.cpp"
"core/runtime/amd_filter_device.cpp"
"core/runtime/amd_topology.cpp"
"core/runtime/default_signal.cpp"
"core/runtime/host_queue.cpp"
"core/runtime/hsa.cpp"
"core/runtime/hsa_api_trace.cpp"
"core/runtime/hsa_ext_amd.cpp"
"core/runtime/hsa_ext_interface.cpp"
"core/runtime/interrupt_signal.cpp"
"core/runtime/intercept_queue.cpp"
"core/runtime/ipc_signal.cpp"
"core/runtime/isa.cpp"
"core/runtime/runtime.cpp"
"core/runtime/signal.cpp"
"core/runtime/queue.cpp"
"core/runtime/cache.cpp"
"core/common/shared.cpp"
"core/common/hsa_table_interface.cpp"
"loader/executable.cpp"
"loader/loaders.cpp"
"libamdhsacode/amd_elf_image.cpp"
"libamdhsacode/amd_hsa_code_util.cpp"
"libamdhsacode/amd_hsa_locks.cpp"
"libamdhsacode/amd_options.cpp"
"libamdhsacode/amd_hsa_code.cpp"
)
set ( SRCS core/util/lnx/os_linux.cpp
core/util/small_heap.cpp
core/util/timer.cpp
core/runtime/amd_blit_kernel.cpp
core/runtime/amd_blit_sdma.cpp
core/runtime/amd_cpu_agent.cpp
core/runtime/amd_gpu_agent.cpp
core/runtime/amd_hsa_loader.cpp
core/runtime/amd_aql_queue.cpp
core/runtime/amd_loader_context.cpp
core/runtime/hsa_ven_amd_loader.cpp
core/runtime/amd_memory_region.cpp
core/runtime/amd_filter_device.cpp
core/runtime/amd_topology.cpp
core/runtime/default_signal.cpp
core/runtime/host_queue.cpp
core/runtime/hsa.cpp
core/runtime/hsa_api_trace.cpp
core/runtime/hsa_ext_amd.cpp
core/runtime/hsa_ext_interface.cpp
core/runtime/interrupt_signal.cpp
core/runtime/intercept_queue.cpp
core/runtime/ipc_signal.cpp
core/runtime/isa.cpp
core/runtime/runtime.cpp
core/runtime/signal.cpp
core/runtime/queue.cpp
core/runtime/cache.cpp
core/common/shared.cpp
core/common/hsa_table_interface.cpp
loader/executable.cpp
loader/loaders.cpp
libamdhsacode/amd_elf_image.cpp
libamdhsacode/amd_hsa_code_util.cpp
libamdhsacode/amd_hsa_locks.cpp
libamdhsacode/amd_options.cpp
libamdhsacode/amd_hsa_code.cpp )
add_library( ${CORE_RUNTIME_TARGET} ${SRCS} )
target_sources( ${CORE_RUNTIME_TARGET} PRIVATE ${SRCS} )
## Build Blit kernels
set(IMAGE_SOURCE_DIR ${CMAKE_CURRENT_SOURCE_DIR}/image)
add_subdirectory(${IMAGE_SOURCE_DIR} ${CMAKE_CURRENT_BINARY_DIR}/image)
if ( NOT DEFINED IMAGE_SUPPORT )
set ( IMAGE_SUPPORT ON )
endif()
set ( IMAGE_SUPPORT ${IMAGE_SUPPORT} CACHE BOOL "Build with image support (default: ON)." )
## Link Core runtime with Thunk (ROCt) and Image libraries
#target_link_libraries ( ${CORE_RUNTIME_TARGET} PUBLIC ${IMAGE_SOURCE_DIR} )
target_link_libraries ( ${CORE_RUNTIME_TARGET} PRIVATE ${IMAGE_TARGET} )
#target_link_libraries ( ${CORE_RUNTIME_TARGET} ${IMAGE_TARGET} )
target_link_libraries ( ${CORE_RUNTIME_TARGET}
PRIVATE hsakmt
elf dl pthread rt
)
## Optional image module defintions.
if(${IMAGE_SUPPORT})
target_compile_definitions( ${CORE_RUNTIME_TARGET} PRIVATE HSA_IMAGE_SUPPORT )
## Strip should be optional or not at all
if ( "${CMAKE_BUILD_TYPE}" STREQUAL Release )
# add_custom_command ( TARGET ${CORE_RUNTIME_TARGET} POST_BUILD COMMAND ${CMAKE_STRIP} $<TARGET_FILE_NAME:${CORE_RUNTIME_TARGET}> )
endif ()
set ( IMAGE_SRCS image/addrlib/src/addrinterface.cpp
image/addrlib/src/core/coord.cpp
image/addrlib/src/core/addrlib.cpp
image/addrlib/src/core/addrlib1.cpp
image/addrlib/src/core/addrlib2.cpp
image/addrlib/src/core/addrobject.cpp
image/addrlib/src/core/addrelemlib.cpp
image/addrlib/src/r800/ciaddrlib.cpp
image/addrlib/src/r800/egbaddrlib.cpp
image/addrlib/src/r800/siaddrlib.cpp
image/addrlib/src/gfx9/gfx9addrlib.cpp
image/addrlib/src/gfx10/gfx10addrlib.cpp
image/device_info.cpp
image/hsa_ext_image.cpp
image/image_runtime.cpp
image/image_manager.cpp
image/image_manager_kv.cpp
image/image_manager_ai.cpp
image/image_manager_nv.cpp
image/image_lut_kv.cpp
image/blit_object_gfx7xx.cpp
image/blit_object_gfx8xx.cpp
image/blit_object_gfx9xx.cpp
image/blit_kernel.cpp
${CMAKE_CURRENT_BINARY_DIR}/image/blit_src/opencl_blit_objects.cpp )
set_source_files_properties(${CMAKE_CURRENT_BINARY_DIR}/image/blit_src/opencl_blit_objects.cpp PROPERTIES GENERATED TRUE)
target_include_directories( ${CORE_RUNTIME_TARGET}
PRIVATE
${CMAKE_CURRENT_SOURCE_DIR}/image
${CMAKE_CURRENT_SOURCE_DIR}/image/addrlib/
${CMAKE_CURRENT_SOURCE_DIR}/image/addrlib/inc
${CMAKE_CURRENT_SOURCE_DIR}/image/addrlib/src
${CMAKE_CURRENT_SOURCE_DIR}/image/addrlib/src/core
${CMAKE_CURRENT_SOURCE_DIR}/image/addrlib/src/r800
${CMAKE_CURRENT_SOURCE_DIR}/image/addrlib/src/gfx9
${CMAKE_CURRENT_SOURCE_DIR}/image/addrlib/src/gfx10
${CMAKE_CURRENT_SOURCE_DIR}/image/addrlib/src/chip/r800
${CMAKE_CURRENT_SOURCE_DIR}/image/addrlib/src/chip/gfx9
${CMAKE_CURRENT_SOURCE_DIR}/image/addrlib/src/chip/gfx10 )
target_sources( ${CORE_RUNTIME_TARGET} PRIVATE ${IMAGE_SRCS} )
## Depend on blit kernel target.
add_subdirectory( ${CMAKE_CURRENT_SOURCE_DIR}/image/blit_src )
add_dependencies( ${CORE_RUNTIME_TARGET} opencl_blit_objects.cpp )
endif()
## Link dependencies.
target_link_libraries ( ${CORE_RUNTIME_TARGET} PRIVATE hsakmt::hsakmt )
target_link_libraries ( ${CORE_RUNTIME_TARGET} PRIVATE elf dl pthread rt )
## Set the VERSION and SOVERSION values
set_property ( TARGET ${CORE_RUNTIME_TARGET} PROPERTY VERSION "${SO_VERSION_STRING}" )
set_property ( TARGET ${CORE_RUNTIME_TARGET} PROPERTY SOVERSION "${SO_MAJOR}" )
set_property ( TARGET ${CORE_RUNTIME_TARGET} PROPERTY SOVERSION "${VERSION_MAJOR}" )
## Create symlinks for legacy packaging and install
add_custom_target ( hsa_include_link ALL WORKING_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR} COMMAND ${CMAKE_COMMAND} -E create_symlink ../hsa/include/hsa hsa_include_link )
add_custom_target ( hsa_include_link ALL WORKING_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR} COMMAND ${CMAKE_COMMAND} -E create_symlink ../../include/hsa hsa_include_link )
if ( ${BUILD_SHARED_LIBS} )
add_custom_target ( hsa_lib_link ALL WORKING_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR} COMMAND ${CMAKE_COMMAND} -E create_symlink ../hsa/lib/${CORE_RUNTIME_LIBRARY}.so ${CORE_RUNTIME_LIBRARY}-link.so )
add_custom_target ( hsa_lib_link2 ALL WORKING_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR} COMMAND ${CMAKE_COMMAND} -E create_symlink ../hsa/lib/${CORE_RUNTIME_LIBRARY}.so.${SO_MAJOR} ${CORE_RUNTIME_LIBRARY}-link.so.${SO_MAJOR} )
add_custom_target ( hsa_lib_link ALL WORKING_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR} COMMAND ${CMAKE_COMMAND} -E create_symlink ../../lib/${CORE_RUNTIME_LIBRARY}.so ${CORE_RUNTIME_LIBRARY}-link.so )
add_custom_target ( hsa_lib_link2 ALL WORKING_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR} COMMAND ${CMAKE_COMMAND} -E create_symlink ../../lib/${CORE_RUNTIME_LIBRARY}.so.${VERSION_MAJOR} ${CORE_RUNTIME_LIBRARY}-link.so.${VERSION_MAJOR} )
endif()
## Set install information
install ( TARGETS ${CORE_RUNTIME_TARGET} DESTINATION hsa/lib COMPONENT binary)
install ( DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/inc/ DESTINATION hsa/include/hsa COMPONENT binary)
install ( FILES ${CMAKE_CURRENT_BINARY_DIR}/hsa_include_link DESTINATION include PERMISSIONS OWNER_WRITE OWNER_READ RENAME hsa COMPONENT dirlink)
# Installs binaries and exports the library usage data to ${HSAKMT_TARGET}Targets
# TODO: Fix me for flat directory layout. Should be ${CMAKE_INSTALL_LIBDIR}
install ( TARGETS ${CORE_RUNTIME_TARGET} EXPORT ${CORE_RUNTIME_TARGET}Targets
ARCHIVE DESTINATION lib COMPONENT binary
LIBRARY DESTINATION lib COMPONENT binary )
# Install license
#install ( FILES ${CMAKE_CURRENT_SOURCE_DIR}/LICENSE.md DESTINATION ${CMAKE_INSTALL_DOCDIR} COMPONENT binary )
# Install public headers
# TODO: Fix me for flat directory layout. Should be ${CMAKE_INSTALL_INCLUDEDIR}
install ( DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/inc/ DESTINATION include/hsa COMPONENT dev )
# Legacy symlink.
install ( FILES ${CMAKE_CURRENT_BINARY_DIR}/hsa_include_link DESTINATION hsa/include PERMISSIONS OWNER_WRITE OWNER_READ RENAME hsa COMPONENT dirlink)
# Legacy symlinks.
if ( ${BUILD_SHARED_LIBS} )
install ( FILES ${CMAKE_CURRENT_BINARY_DIR}/${CORE_RUNTIME_LIBRARY}-link.so DESTINATION lib PERMISSIONS OWNER_WRITE OWNER_READ RENAME ${CORE_RUNTIME_LIBRARY}.so COMPONENT binary)
install ( FILES ${CMAKE_CURRENT_BINARY_DIR}/${CORE_RUNTIME_LIBRARY}-link.so.${SO_MAJOR} DESTINATION lib PERMISSIONS OWNER_WRITE OWNER_READ RENAME ${CORE_RUNTIME_LIBRARY}.so.${SO_MAJOR} COMPONENT binary)
install ( FILES ${CMAKE_CURRENT_BINARY_DIR}/${CORE_RUNTIME_LIBRARY}-link.so DESTINATION hsa/lib PERMISSIONS OWNER_WRITE OWNER_READ RENAME ${CORE_RUNTIME_LIBRARY}.so COMPONENT binary)
install ( FILES ${CMAKE_CURRENT_BINARY_DIR}/${CORE_RUNTIME_LIBRARY}-link.so.${VERSION_MAJOR} DESTINATION hsa/lib PERMISSIONS OWNER_WRITE OWNER_READ RENAME ${CORE_RUNTIME_LIBRARY}.so.${VERSION_MAJOR} COMPONENT binary)
endif ()
## Configure and install package config file
# Record our usage data for clients find_package calls.
# TODO: Fix me for flat directory layout. Should be ${CMAKE_INSTALL_LIBDIR}
install ( EXPORT ${CORE_RUNTIME_TARGET}Targets
FILE ${CORE_RUNTIME_TARGET}Targets.cmake
NAMESPACE ${CORE_RUNTIME_TARGET}::
DESTINATION lib/cmake/${CORE_RUNTIME_TARGET}
COMPONENT dev)
# Adds the target alias hsa-runtime64::hsa-runtime64 to the local cmake cache.
# This isn't necessary today. It's harmless preparation for some
# hypothetical future in which the we might be inluded by add_subdirectory()
# in some other project's cmake file. It allows uniform use of find_package
# and target_link_library() without regard to whether a target is external or
# a subdirectory of the current build.
add_library( ${CORE_RUNTIME_TARGET}::${CORE_RUNTIME_TARGET} ALIAS ${CORE_RUNTIME_TARGET} )
# Create cmake configuration files
include(CMakePackageConfigHelpers)
# TODO: Fix me for flat directory layout. Should be ${CMAKE_INSTALL_LIBDIR}
configure_package_config_file(${CORE_RUNTIME_TARGET}-config.cmake.in
${CORE_RUNTIME_TARGET}-config.cmake
INSTALL_DESTINATION lib/cmake/${CORE_RUNTIME_TARGET} )
write_basic_package_version_file(${CORE_RUNTIME_TARGET}-config-version.cmake
VERSION ${SO_VERSION_STRING} COMPATIBILITY AnyNewerVersion )
# TODO: Fix me for flat directory layout. Should be ${CMAKE_INSTALL_LIBDIR}
install(FILES ${CMAKE_CURRENT_BINARY_DIR}/${CORE_RUNTIME_TARGET}-config.cmake ${CMAKE_CURRENT_BINARY_DIR}/${CORE_RUNTIME_TARGET}-config-version.cmake
DESTINATION lib/cmake/${CORE_RUNTIME_TARGET}
COMPONENT dev)
## Packaging directives
set ( CPACK_GENERATOR "DEB;RPM" CACHE STRING "Package types to build")
## Only pack the "binary" components, post install script will add the directory link.
## Only pack the "binary" and "dev" components, post install script will add the directory link.
set (CPACK_DEB_COMPONENT_INSTALL ON)
set (CPACK_COMPONENTS_ALL_IN_ONE_PACKAGE 1)
set (CPACK_COMPONENTS_ALL binary)
set (CPACK_COMPONENTS_ALL binary dev)
set ( CPACK_PACKAGE_NAME "hsa-rocr-dev" )
set ( CPACK_PACKAGE_VENDOR "AMD" )
@@ -234,13 +325,13 @@ set ( CPACK_PACKAGE_DESCRIPTION_SUMMARY "AMD Heterogeneous System Architecture H
set ( CPACK_RESOURCE_FILE_LICENSE "${CMAKE_CURRENT_SOURCE_DIR}/LICENSE.md" )
## Process the install scripts to update the CPACK variables
configure_file(${CMAKE_CURRENT_SOURCE_DIR}/DEBIAN/post_install ${CMAKE_CURRENT_SOURCE_DIR}/DEBIAN/postinst @ONLY)
configure_file(${CMAKE_CURRENT_SOURCE_DIR}/DEBIAN/pre_remove ${CMAKE_CURRENT_SOURCE_DIR}/DEBIAN/prerm @ONLY)
configure_file(${CMAKE_CURRENT_SOURCE_DIR}/DEBIAN/post_install DEBIAN/postinst @ONLY)
configure_file(${CMAKE_CURRENT_SOURCE_DIR}/DEBIAN/pre_remove DEBIAN/prerm @ONLY)
# Debian package specific variables
set ( CPACK_DEBIAN_PACKAGE_DEPENDS "hsakmt-roct" )
set ( CPACK_DEBIAN_PACKAGE_HOMEPAGE "https://github.com/RadeonOpenCompute/ROCR-Runtime" )
set ( CPACK_DEBIAN_PACKAGE_CONTROL_EXTRA "${CMAKE_CURRENT_SOURCE_DIR}/DEBIAN/postinst;${CMAKE_CURRENT_SOURCE_DIR}/DEBIAN/prerm" )
set ( CPACK_DEBIAN_PACKAGE_CONTROL_EXTRA "DEBIAN/postinst;DEBIAN/prerm" )
# Declare that this package will replace functionality provided by hsa-ext-rocr-dev package
set ( CPACK_DEBIAN_PACKAGE_BREAKS "hsa-ext-rocr-dev" )
@@ -252,7 +343,6 @@ set ( CPACK_DEBIAN_PACKAGE_CONFLICTS "hsa-ext-rocr-dev" )
set ( CPACK_RPM_PACKAGE_DEPENDS "hsakmt-roct" )
# Declare that this package will replace functionality provided by hsa-ext-rocr-dev package
#set ( CPACK_DEBIAN_PACKAGE_BREAKS "hsa-ext-rocr-dev" ) # equivalent command not found
set ( CPACK_RPM_PACKAGE_PROVIDES "hsa-ext-rocr-dev" )
set ( CPACK_RPM_PACKAGE_OBSOLETES "hsa-ext-rocr-dev" )
set ( CPACK_RPM_PACKAGE_CONFLICTS "hsa-ext-rocr-dev" )
projects/rocr-runtime/runtime/hsa-runtime/DEBIAN/post_install
+4 -3
Vedi File
@@ -1,4 +1,4 @@
#/bin/bash
#!/bin/bash
set -e
@@ -9,8 +9,9 @@ do_ldconfig() {
case "$1" in
configure)
do_ldconfig
mkdir -p @CPACK_PACKAGING_INSTALL_PREFIX@/include
ln -sf ../hsa/include/hsa @CPACK_PACKAGING_INSTALL_PREFIX@/include/hsa
# Workaround for CPACK directory symlink handling error.
mkdir -p @CPACK_PACKAGING_INSTALL_PREFIX@/hsa/include
ln -sf ../../include/hsa @CPACK_PACKAGING_INSTALL_PREFIX@/hsa/include/hsa
;;
abort-upgrade|abort-remove|abort-deconfigure)
echo "$1"
projects/rocr-runtime/runtime/hsa-runtime/DEBIAN/pre_remove
+2 -1
Vedi File
@@ -9,7 +9,8 @@ rm_ldconfig() {
case "$1" in
remove)
rm_ldconfig
rm -rf @CPACK_PACKAGING_INSTALL_PREFIX@/include/hsa
# Workaround for CPACK directory symlink handling error.
rm -rf @CPACK_PACKAGING_INSTALL_PREFIX@/hsa/include/hsa
;;
purge)
;;
projects/rocr-runtime/runtime/hsa-runtime/cmake_modules/hsa_common.cmake
+1 -1
Vedi File
@@ -53,7 +53,7 @@ endif()
if(UNIX)
set(PS ":")
set(CMAKE_CXX_FLAGS "-Wall -std=c++11 ${EXTRA_CFLAGS}")
set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -Wall -std=c++11 ${EXTRA_CFLAGS}")
set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -fpic")
if (CMAKE_COMPILER_IS_GNUCXX)
set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -Wl,--unresolved-symbols=ignore-in-shared-libs")
projects/rocr-runtime/runtime/hsa-runtime/core/runtime/hsa_ext_interface.cpp
+4
Vedi File
@@ -122,7 +122,9 @@ void ExtensionEntryPoints::UnloadImage() {
InitAmdExtTable();
InitImageExtTable();
core::hsa_internal_api_table_.Reset();
#ifdef HSA_IMAGE_SUPPORT
rocr::image::ReleaseImageRsrcs();
#endif
}
void ExtensionEntryPoints::Unload() {
@@ -151,6 +153,7 @@ void ExtensionEntryPoints::Unload() {
}
bool ExtensionEntryPoints::LoadImage() {
#ifdef HSA_IMAGE_SUPPORT
// Consult user input on linking to Image implementation
bool disable_image = core::Runtime::runtime_singleton_->flag().disable_image();
if (disable_image) {
@@ -172,6 +175,7 @@ bool ExtensionEntryPoints::LoadImage() {
// Update Amd Ext Api table Api that deals with Images
UpdateAmdExtTable(func);
#endif
return true;
}
projects/rocr-runtime/runtime/hsa-runtime/hsa-runtime64-config.cmake.in
+7
Vedi File
@@ -0,0 +1,7 @@
@PACKAGE_INIT@
include( CMakeFindDependencyMacro )
find_dependency(hsakmt 1.0)
include( "${CMAKE_CURRENT_LIST_DIR}/@CORE_RUNTIME_TARGET@Targets.cmake" )
projects/rocr-runtime/runtime/hsa-runtime/image/CMakeLists.txt
-182
Vedi File
@@ -1,182 +0,0 @@
cmake_minimum_required ( VERSION 3.5.0 )
## Set ext runtime module name and project name.
set ( IMAGE_NAME "hsa-ext-image" )
set ( IMAGE_TARGET "${IMAGE_NAME}64" )
set ( IMAGE_LIBRARY "lib${IMAGE_TARGET}" )
project ( ${IMAGE_TARGET} )
## Build image as a shared or static library
#if ( NOT DEFINED BUILD_SHARED_LIBS )
# set ( BUILD_SHARED_LIBS "on" )
#endif()
#set ( BUILD_SHARED_LIBS ${BUILD_SHARED_LIBS} CACHE BOOL "Build shared library (.so) or not.")
# Optionally, build with ccache.
set(ROCM_CCACHE_BUILD OFF CACHE BOOL "Set to ON for a ccache enabled build")
if (ROCM_CCACHE_BUILD)
find_program(CCACHE_PROGRAM ccache)
if (CCACHE_PROGRAM)
set_property(GLOBAL PROPERTY RULE_LAUNCH_COMPILE ${CCACHE_PROGRAM})
else()
message(WARNING "Unable to find ccache. Falling back to real compiler")
endif() # if (CCACHE_PROGRAM)
endif() # if (ROCM_CCACHE_BUILD)
## Include the cmake_modules utils.cmake
list ( APPEND CMAKE_MODULE_PATH "${CMAKE_CURRENT_SOURCE_DIR}/../cmake_modules" )
include ( utils )
## Compiler preproc definitions.
add_definitions ( -D__linux__ )
add_definitions ( -DUNIX_OS )
add_definitions ( -DLINUX )
add_definitions ( -D__AMD64__ )
add_definitions ( -D__x86_64__ )
add_definitions ( -DAMD_INTERNAL_BUILD )
add_definitions ( -DLITTLEENDIAN_CPU=1 )
add_definitions ( -D HSA_DEPRECATED= )
add_definitions ( -D BRAHMA_BUILD=1 )
## Get the package version.
#get_version( "1.1.9")
#set(SO_MAJOR 1)
#set(SO_MINOR 1)
#if ( ${ROCM_PATCH_VERSION} )
# set ( SO_PATCH ${ROCM_PATCH_VERSION})
# set ( VERSION_PATCH ${ROCM_PATCH_VERSION})
#else ()
# set(SO_PATCH 9)
#endif ()
#set( SO_VERSION_STRING "${SO_MAJOR}.${SO_MINOR}.${SO_PATCH}" )
#set( PACKAGE_VERSION_STRING "${VERSION_MAJOR}.${VERSION_MINOR}.${VERSION_PATCH}.${VERSION_COMMIT_COUNT}-${VERSION_JOB}-${VERSION_HASH}" )
## Find the hsakmt library and include files
get_include_path( HSAKMT_INC_PATH "libhsakmt include path" NAMES "hsakmt.h" "libhsakmt/hsakmt.h" HINTS "${CMAKE_BINARY_DIR}/../../include" "${CMAKE_CURRENT_SOURCE_DIR}/../../../../../libhsakmt/include" PATHS "/opt/rocm/include")
get_library_path( HSAKMT_LIB_PATH "libhsakmt library path" NAMES "libhsakmt.so" HINTS "${CMAKE_BINARY_DIR}/../../lib" "${CMAKE_BINARY_DIR}/../roct" PATHS "/opt/rocm/lib")
include_directories( ${HSAKMT_INC_PATH} )
link_directories( ${HSAKMT_LIB_PATH} )
## Find the hsa-runtime and include files
#get_include_path( HSA_INC_PATH "ROCr include path" NAMES "inc/hsa.h" HINTS "${CMAKE_CURRENT_SOURCE_DIR}/../hsa-runtime" "${CMAKE_BINARY_DIR}/../../include" PATHS "/opt/rocm/include")
#get_library_path( HSA_LIB_PATH "ROCr library path" NAMES "libhsa-runtime64.so" HINTS "${CMAKE_BINARY_DIR}/../../lib" "${CMAKE_BINARY_DIR}/../hsa-core" "${CMAKE_CURRENT_SOURCE_DIR}/../hsa-runtime/build" PATHS "/opt/rocm/lib")
#include_directories( ${HSA_INC_PATH} )
#link_directories( ${HSA_LIB_PATH} )
## Find self
if( "${IMAGE_SOURCE_DIR}" STREQUAL "" )
get_include_path( IMG_SOURCE_FILE null NAMES "image_runtime.h" HINTS "${CMAKE_CURRENT_SOURCE_DIR}" )
get_filename_component( IMAGE_SOURCE_DIR "${IMG_SOURCE_FILE}" ABSOLUTE )
unset( IMG_SOURCE_FILE CACHE )
endif()
set( IMAGE_SOURCE_DIR ${IMAGE_SOURCE_DIR} CACHE PATH "Image lib source dir" FORCE )
message( "Value of cmake_current_source_dir: ${CMAKE_CURRENT_SOURCE_DIR} ")
get_filename_component( OPEN_SOURCE_DIR "${CMAKE_CURRENT_SOURCE_DIR}/../.." ABSOLUTE CACHE )
message( "Value of open_source_dir: ${OPEN_SOURCE_DIR} ")
set( OPEN_SOURCE_DIR ${OPEN_SOURCE_DIR} CACHE PATH "Open source root dir" FORCE )
## Set RUNPATH - ../../lib covers use of the legacy symlink in /hsa/lib/
set(CMAKE_INSTALL_RPATH "$ORIGIN;$ORIGIN/../../lib;$ORIGIN/../../lib64;$ORIGIN/../lib64")
## ------------------------- Linux Compiler and Linker options -------------------------
set ( CMAKE_CXX_FLAGS "-std=c++11 " )
#set ( CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -Wno-deprecated-declarations -fPIC" )
#set ( CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -fPIC" )
set ( CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -Werror -fexceptions -fno-rtti -fvisibility=hidden -Wno-error=sign-compare -Wno-sign-compare -Wno-write-strings -Wno-deprecated-declarations -Wno-conversion-null -fno-math-errno -fno-threadsafe-statics -fmerge-all-constants -fms-extensions -Wno-error=comment -Wno-comment -Wno-error=pointer-arith -Wno-pointer-arith -fPIC" )
#if ( CMAKE_SYSTEM_PROCESSOR STREQUAL "x86_64" )
# set ( CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -m64 -msse -msse2" )
#elseif ( CMAKE_SYSTEM_PROCESSOR STREQUAL "x86" )
# set ( CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -m32" )
#endif ()
#if ( "${CMAKE_BUILD_TYPE}" STREQUAL Debug )
# set( CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -ggdb" )
#endif ()
#set ( DRVDEF "${IMAGE_SOURCE_DIR}/image.so.def" )
set ( CMAKE_SHARED_LINKER_FLAGS "-Wl,-Bdynamic -Wl,-z,noexecstack -Wl,--version-script=${DRVDEF} -Wl,--enable-new-dtags" )
## Library path(s).
include_directories(${IMAGE_SOURCE_DIR}/)
include_directories(${OPEN_SOURCE_DIR}/hsa-runtime)
include_directories(${OPEN_SOURCE_DIR}/hsa-runtime/inc)
include_directories(${OPEN_SOURCE_DIR}/hsa-runtime/core/inc)
include_directories(${IMAGE_SOURCE_DIR}/addrlib/)
include_directories(${IMAGE_SOURCE_DIR}/addrlib/inc)
include_directories(${IMAGE_SOURCE_DIR}/addrlib/src)
include_directories(${IMAGE_SOURCE_DIR}/addrlib/src/core)
include_directories(${IMAGE_SOURCE_DIR}/addrlib/src/r800)
include_directories(${IMAGE_SOURCE_DIR}/addrlib/src/gfx9)
include_directories(${IMAGE_SOURCE_DIR}/addrlib/src/gfx10)
include_directories(${IMAGE_SOURCE_DIR}/addrlib/src/chip/r800)
include_directories(${IMAGE_SOURCE_DIR}/addrlib/src/chip/gfx9)
include_directories(${IMAGE_SOURCE_DIR}/addrlib/src/chip/gfx10)
set ( IMAGE_SRCS ${IMAGE_SOURCE_DIR}/addrlib/src/addrinterface.cpp
${IMAGE_SOURCE_DIR}/addrlib/src/core/coord.cpp
${IMAGE_SOURCE_DIR}/addrlib/src/core/addrlib.cpp
${IMAGE_SOURCE_DIR}/addrlib/src/core/addrlib1.cpp
${IMAGE_SOURCE_DIR}/addrlib/src/core/addrlib2.cpp
${IMAGE_SOURCE_DIR}/addrlib/src/core/addrobject.cpp
${IMAGE_SOURCE_DIR}/addrlib/src/core/addrelemlib.cpp
${IMAGE_SOURCE_DIR}/addrlib/src/r800/ciaddrlib.cpp
${IMAGE_SOURCE_DIR}/addrlib/src/r800/egbaddrlib.cpp
${IMAGE_SOURCE_DIR}/addrlib/src/r800/siaddrlib.cpp
${IMAGE_SOURCE_DIR}/addrlib/src/gfx9/gfx9addrlib.cpp
${IMAGE_SOURCE_DIR}/addrlib/src/gfx10/gfx10addrlib.cpp
${IMAGE_SOURCE_DIR}/device_info.cpp
${IMAGE_SOURCE_DIR}/hsa_ext_image.cpp
${IMAGE_SOURCE_DIR}/image_runtime.cpp
${IMAGE_SOURCE_DIR}/image_manager.cpp
${IMAGE_SOURCE_DIR}/image_manager_kv.cpp
${IMAGE_SOURCE_DIR}/image_manager_ai.cpp
${IMAGE_SOURCE_DIR}/image_manager_nv.cpp
${IMAGE_SOURCE_DIR}/image_lut_kv.cpp
${IMAGE_SOURCE_DIR}/blit_object_gfx7xx.cpp
${IMAGE_SOURCE_DIR}/blit_object_gfx8xx.cpp
${IMAGE_SOURCE_DIR}/blit_object_gfx9xx.cpp
${IMAGE_SOURCE_DIR}/opencl_blit_objects.cpp
${IMAGE_SOURCE_DIR}/blit_kernel.cpp
)
add_subdirectory(${IMAGE_SOURCE_DIR}/blit_src ${CMAKE_CURRENT_BINARY_DIR}/image_blit)
set_source_files_properties(${IMAGE_SOURCE_DIR}/opencl_blit_objects.cpp PROPERTIES GENERATED 1)
#add_library ( ${IMAGE_TARGET} SHARED ${IMAGE_SRCS} )
add_library ( ${IMAGE_TARGET} STATIC ${IMAGE_SRCS} )
add_dependencies( ${IMAGE_TARGET} opencl_blit_objects.cpp )
## Set the VERSION and SOVERSION values
#set_property ( TARGET ${IMAGE_TARGET} PROPERTY VERSION "${SO_VERSION_STRING}" )
#set_property ( TARGET ${IMAGE_TARGET} PROPERTY SOVERSION "${SO_MAJOR}" )
## Add the core runtime in the link
#target_link_libraries (
# ${IMAGE_TARGET}
# PRIVATE hsa-runtime64
# PRIVATE hsakmt
# c dl pthread rt
#)
## If the build is Release, strip the target library
#if ( "${CMAKE_BUILD_TYPE}" STREQUAL Release )
# add_custom_command ( TARGET ${IMAGE_TARGET} POST_BUILD COMMAND ${CMAKE_STRIP} $<TARGET_FILE_NAME:${IMAGE_TARGET}> )
#endif ()
## Create symlinks for legacy packaging and install
#if ( ${BUILD_SHARED_LIBS} )
# add_custom_target ( hsa_images_lib_link ALL WORKING_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR} COMMAND ${CMAKE_COMMAND} -E create_symlink ../hsa/lib/${IMAGE_LIBRARY}.so ${IMAGE_LIBRARY}-link.so )
#endif()
## Set install information
#if ( ${BUILD_SHARED_LIBS} )
# install ( TARGETS ${IMAGE_TARGET} LIBRARY DESTINATION hsa/lib )
# install ( FILES ${CMAKE_CURRENT_BINARY_DIR}/${IMAGE_LIBRARY}-link.so DESTINATION lib PERMISSIONS OWNER_WRITE OWNER_READ RENAME ${IMAGE_LIBRARY}.so )
#endif()
projects/rocr-runtime/runtime/hsa-runtime/image/blit_kernel_cl.cl
-448
Vedi File
@@ -1,448 +0,0 @@
/// Kernel code for HSA image import/export/copy/clear in OpenCL C form.
__kernel void copy_image_to_buffer(
__read_only image2d_array_t src,
__global uint* dstUInt,
__global ushort* dstUShort,
__global uchar* dstUChar,
int4 srcOrigin,
ulong4 dstOrigin,
int4 size,
uint4 format,
ulong4 pitch)
{
ulong idxDst;
int4 coordsSrc;
uint4 texel;
coordsSrc.x = get_global_id(0);
coordsSrc.y = get_global_id(1);
coordsSrc.z = get_global_id(2);
coordsSrc.w = 0;
if ((coordsSrc.x >= size.x) ||
(coordsSrc.y >= size.y) ||
(coordsSrc.z >= size.z)) {
return;
}
idxDst = (coordsSrc.z * pitch.y + coordsSrc.y * pitch.x +
coordsSrc.x) * format.z + dstOrigin.x;
coordsSrc.x += srcOrigin.x;
coordsSrc.y += srcOrigin.y;
coordsSrc.z += srcOrigin.z;
texel = read_imageui(src, coordsSrc);
// Check components
switch (format.x) {
case 1:
// Check size
switch (format.y) {
case 1:
dstUChar[idxDst] = (uchar)texel.x;
break;
case 2:
dstUShort[idxDst] = (ushort)texel.x;
break;
case 4:
dstUInt[idxDst] = texel.x;
break;
}
break;
case 2:
// Check size
switch (format.y) {
case 1:
dstUShort[idxDst] = (ushort)texel.x |
((ushort)texel.y << 8);
break;
case 2:
dstUInt[idxDst] = texel.x | (texel.y << 16);
break;
case 4:
dstUInt[idxDst++] = texel.x;
dstUInt[idxDst] = texel.y;
break;
}
break;
case 4:
// Check size
switch (format.y) {
case 1:
dstUInt[idxDst] = (uint)texel.x |
(texel.y << 8) |
(texel.z << 16) |
(texel.w << 24);
break;
case 2:
dstUInt[idxDst++] = texel.x | (texel.y << 16);
dstUInt[idxDst] = texel.z | (texel.w << 16);
break;
case 4:
dstUInt[idxDst++] = texel.x;
dstUInt[idxDst++] = texel.y;
dstUInt[idxDst++] = texel.z;
dstUInt[idxDst] = texel.w;
break;
}
break;
}
}
__kernel void copy_buffer_to_image(
__global uint* src,
__write_only image2d_array_t dst,
ulong4 srcOrigin,
int4 dstOrigin,
int4 size,
uint4 format,
ulong4 pitch)
{
ulong idxSrc;
int4 coordsDst;
uint4 pixel;
__global uint* srcUInt = src;
__global ushort* srcUShort = (__global ushort*)src;
__global uchar* srcUChar = (__global uchar*)src;
ushort tmpUShort;
uint tmpUInt;
coordsDst.x = get_global_id(0);
coordsDst.y = get_global_id(1);
coordsDst.z = get_global_id(2);
coordsDst.w = 0;
if ((coordsDst.x >= size.x) ||
(coordsDst.y >= size.y) ||
(coordsDst.z >= size.z)) {
return;
}
idxSrc = (coordsDst.z * pitch.y +
coordsDst.y * pitch.x + coordsDst.x) *
format.z + srcOrigin.x;
coordsDst.x += dstOrigin.x;
coordsDst.y += dstOrigin.y;
coordsDst.z += dstOrigin.z;
// Check components
switch (format.x) {
case 1:
// Check size
if (format.y == 1) {
pixel.x = (uint)srcUChar[idxSrc];
}
else if (format.y == 2) {
pixel.x = (uint)srcUShort[idxSrc];
}
else {
pixel.x = srcUInt[idxSrc];
}
break;
case 2:
// Check size
if (format.y == 1) {
tmpUShort = srcUShort[idxSrc];
pixel.x = (uint)(tmpUShort & 0xff);
pixel.y = (uint)(tmpUShort >> 8);
}
else if (format.y == 2) {
tmpUInt = srcUInt[idxSrc];
pixel.x = (tmpUInt & 0xffff);
pixel.y = (tmpUInt >> 16);
}
else {
pixel.x = srcUInt[idxSrc++];
pixel.y = srcUInt[idxSrc];
}
break;
case 4:
// Check size
if (format.y == 1) {
tmpUInt = srcUInt[idxSrc];
pixel.x = tmpUInt & 0xff;
pixel.y = (tmpUInt >> 8) & 0xff;
pixel.z = (tmpUInt >> 16) & 0xff;
pixel.w = (tmpUInt >> 24) & 0xff;
}
else if (format.y == 2) {
tmpUInt = srcUInt[idxSrc++];
pixel.x = tmpUInt & 0xffff;
pixel.y = (tmpUInt >> 16);
tmpUInt = srcUInt[idxSrc];
pixel.z = tmpUInt & 0xffff;
pixel.w = (tmpUInt >> 16);
}
else {
pixel.x = srcUInt[idxSrc++];
pixel.y = srcUInt[idxSrc++];
pixel.z = srcUInt[idxSrc++];
pixel.w = srcUInt[idxSrc];
}
break;
}
// Write the final pixel
write_imageui(dst, coordsDst, pixel);
}
__kernel void copy_image_default(
__read_only image2d_array_t src,
__write_only image2d_array_t dst,
int4 srcOrigin,
int4 dstOrigin,
int4 size)
{
int4 coordsDst;
int4 coordsSrc;
coordsDst.x = get_global_id(0);
coordsDst.y = get_global_id(1);
coordsDst.z = get_global_id(2);
coordsDst.w = 0;
if ((coordsDst.x >= size.x) ||
(coordsDst.y >= size.y) ||
(coordsDst.z >= size.z)) {
return;
}
coordsSrc = srcOrigin + coordsDst;
coordsDst += dstOrigin;
uint4 texel;
texel = read_imageui(src, coordsSrc);
write_imageui(dst, coordsDst, texel);
}
float linear_to_standard_rgba(float l_val) {
float s_val = l_val;
if (isnan(s_val)) s_val = 0.0f;
if (s_val > 1.0f) {
s_val = 1.0f;
} else if (s_val < 0.0f) {
s_val = 0.0f;
} else if (s_val < 0.0031308f) {
s_val = 12.92f * s_val;
} else {
s_val = (1.055f * pow(s_val, 5.0f / 12.0f)) - 0.055f;
}
return s_val;
}
__kernel void copy_image_linear_to_standard(
__read_only image2d_array_t src,
__write_only image2d_array_t dst,
int4 srcOrigin,
int4 dstOrigin,
int4 size,
int copyType)
{
int4 coordsDst;
int4 coordsSrc;
coordsDst.x = get_global_id(0);
coordsDst.y = get_global_id(1);
coordsDst.z = get_global_id(2);
coordsDst.w = 0;
if ((coordsDst.x >= size.x) ||
(coordsDst.y >= size.y) ||
(coordsDst.z >= size.z)) {
return;
}
coordsSrc = srcOrigin + coordsDst;
coordsDst += dstOrigin;
float4 texel;
texel = read_imagef(src, coordsSrc);
texel.x = linear_to_standard_rgba(texel.x);
texel.y = linear_to_standard_rgba(texel.y);
texel.z = linear_to_standard_rgba(texel.z);
write_imagef(dst, coordsDst, texel);
}
__kernel void copy_image_standard_to_linear(
__read_only image2d_array_t src,
__write_only image2d_array_t dst,
int4 srcOrigin,
int4 dstOrigin,
int4 size,
int copyType)
{
int4 coordsDst;
int4 coordsSrc;
coordsDst.x = get_global_id(0);
coordsDst.y = get_global_id(1);
coordsDst.z = get_global_id(2);
coordsDst.w = 0;
if ((coordsDst.x >= size.x) ||
(coordsDst.y >= size.y) ||
(coordsDst.z >= size.z)) {
return;
}
coordsSrc = srcOrigin + coordsDst;
coordsDst += dstOrigin;
float4 texel;
texel = read_imagef(src, coordsSrc);
write_imagef(dst, coordsDst, texel);
}
__kernel void copy_image_1db(
__read_only image1d_buffer_t src,
__write_only image1d_buffer_t dst,
int4 srcOrigin,
int4 dstOrigin,
int4 size)
{
int coordDst;
int coordSrc;
coordDst = get_global_id(0);
if (coordDst >= size.x) {
return;
}
coordSrc = srcOrigin.x + coordDst;
coordDst += dstOrigin.x;
uint4 texel;
texel = read_imageui(src, coordSrc);
write_imageui(dst, coordDst, texel);
}
__kernel void copy_image_1db_to_reg(
__read_only image1d_buffer_t src,
__write_only image2d_array_t dst,
int4 srcOrigin,
int4 dstOrigin,
int4 size)
{
int4 coordsDst;
int coordSrc;
coordsDst.x = get_global_id(0);
coordsDst.y = get_global_id(1);
coordsDst.z = get_global_id(2);
coordsDst.w = 0;
if (coordsDst.x >= size.x) {
return;
}
coordSrc = srcOrigin.x + coordsDst.x;
coordsDst += dstOrigin;
uint4 texel;
texel = read_imageui(src, coordSrc);
write_imageui(dst, coordsDst, texel);
}
__kernel void copy_image_reg_to_1db(
__read_only image2d_array_t src,
__write_only image1d_buffer_t dst,
int4 srcOrigin,
int4 dstOrigin,
int4 size)
{
int coordDst;
int4 coordsSrc;
coordsSrc.x = get_global_id(0);
coordsSrc.y = get_global_id(1);
coordsSrc.z = get_global_id(2);
coordsSrc.w = 0;
if (coordsSrc.x >= size.x) {
return;
}
coordDst = dstOrigin.x + coordsSrc.x;
coordsSrc += srcOrigin;
uint4 texel;
texel = read_imageui(src, coordsSrc);
write_imageui(dst, coordDst, texel);
}
__kernel void clear_image(
__write_only image2d_array_t image,
float4 patternFLOAT4,
int4 patternINT4,
uint4 patternUINT4,
int4 origin,
int4 size,
uint type)
{
int4 coords;
coords.x = get_global_id(0);
coords.y = get_global_id(1);
coords.z = get_global_id(2);
coords.w = 0;
if ((coords.x >= size.x) ||
(coords.y >= size.y) ||
(coords.z >= size.z)) {
return;
}
coords += origin;
// Check components
switch (type) {
case 0:
write_imagef(image, coords, patternFLOAT4);
break;
case 1:
write_imagei(image, coords, patternINT4);
break;
case 2:
write_imageui(image, coords, patternUINT4);
break;
}
}
__kernel void clear_image_1db(
__write_only image1d_buffer_t image,
float4 patternFLOAT4,
int4 patternINT4,
uint4 patternUINT4,
int4 origin,
int4 size,
uint type)
{
int coord = get_global_id(0);
if (coord >= size.x) {
return;
}
coord += origin.x;
// Check components
switch (type) {
case 0:
write_imagef(image, coord, patternFLOAT4);
break;
case 1:
write_imagei(image, coord, patternINT4);
break;
case 2:
write_imageui(image, coord, patternUINT4);
break;
}
}
projects/rocr-runtime/runtime/hsa-runtime/image/blit_src/CMakeLists.txt
+92 -116
Vedi File
@@ -1,92 +1,61 @@
#
# Minimum version of cmake required
#
cmake_minimum_required(VERSION 3.5.0)
################################################################################
##
## The University of Illinois/NCSA
## Open Source License (NCSA)
##
## Copyright (c) 2014-2017, Advanced Micro Devices, Inc. All rights reserved.
##
## Developed by:
##
## AMD Research and AMD HSA Software Development
##
## Advanced Micro Devices, Inc.
##
## www.amd.com
##
## Permission is hereby granted, free of charge, to any person obtaining a copy
## of this software and associated documentation files (the "Software"), to
## deal with 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:
##
## - Redistributions of source code must retain the above copyright notice,
## this list of conditions and the following disclaimers.
## - Redistributions in binary form must reproduce the above copyright
## notice, this list of conditions and the following disclaimers in
## the documentation and#or other materials provided with the distribution.
## - Neither the names of Advanced Micro Devices, Inc,
## nor the names of its contributors may be used to endorse or promote
## products derived from this Software without specific prior written
## permission.
##
## 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 CONTRIBUTORS 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 WITH THE SOFTWARE.
##
################################################################################
#
# Required Defines on cmake command line
#
# 1) Set location of OpenCL header files
# OPENCL_DIR="Root for OpenCL install"
# If not set, the default value is "/opt/rocm/opencl"
#
# 2) Set location of CLANG/LLVM binary directory
# LLVM_DIR="Directory contains clang, llvm-link and llvm-dis
# If not set, the default value is "<PROJECT_BUILD_DIR>/lightning/bin"
#
# 3) Set BITCODE library directory
# BITCODE_DIR="Directory contains the bitcode library"
# If not set, the default value is "${OPENCL_DIR}/lib/x86_64/bitcode"
#
# 4) Set TARGET_DEVICES to indicate gpu types for kernel builds (e.g., "gfx803;gfx900; ...")
# If not set, the target devices are those have the Open Compute Library Controls (OCLC)
# bitcode file, "oclc_isa_version_*.amdgcn.bc", in the BITCODE directory
#
# Building - Should be automatic but for manual builds:
#
# 1) *** Create build folder e.g. "blit_src/build" - any name will do
# 2) Go to the build folder
# 3) Run "cmake .."
# 4) Run "make opencl_blit_objects.cpp"
#
## Include the cmake_modules utils.cmake
list ( APPEND CMAKE_MODULE_PATH "${CMAKE_CURRENT_SOURCE_DIR}/../../cmake_modules" )
include ( utils )
cmake_minimum_required ( VERSION 3.5.0 )
# Flag to abort before executing after default initialization of cache variables
set (QUIT 0)
# Collect possible LLVM version directories.
set (LLVM_SEARCH_PATHS "")
set (LLVM_SEARCH_ROOT "${CMAKE_INSTALL_PREFIX}/llvm/lib/clang")
listsubdirs(${LLVM_SEARCH_ROOT} FOLDERS)
foreach(ITEM IN LISTS FOLDERS)
list (APPEND LLVM_SEARCH_PATHS "${LLVM_SEARCH_ROOT}/${ITEM}/include/")
endforeach()
# Import target 'clang'
find_package(Clang REQUIRED HINTS ${CMAKE_INSTALL_PREFIX}/llvm PATHS /opt/rocm/llvm )
if (NOT DEFINED OPENCL_VER)
set (OPENCL_VER "2.0")
endif()
set( OPENCL_VER ${OPENCL_VER} CACHE STRING "OpenCL version" FORCE )
get_include_path(BITCODE_DIR "Bitcode library path" RESULT FOUND NAMES "opencl.amdgcn.bc" HINTS "${CMAKE_INSTALL_PREFIX}/lib/bitcode" "${OPENCL_DIR}/lib/x86_64/bitcode")
# Device libs doesn't support find_package yet.
get_include_path(BITCODE_DIR "Bitcode library path" RESULT FOUND NAMES "opencl.amdgcn.bc"
HINTS "${CMAKE_INSTALL_PREFIX}/lib/bitcode" "${CMAKE_INSTALL_PREFIX}/lib/x86_64/bitcode"
PATHS "/opt/rocm/lib/bitcode" "${CMAKE_INSTALL_PREFIX}/lib/x86_64/bitcode")
if (NOT ${FOUND})
set (QUIT 1)
endif()
set (BITCODE_LIB "${BITCODE_DIR}/opencl.amdgcn.bc")
if (NOT EXISTS ${BITCODE_LIB})
message("ERROR: path to opencl.amdgcn.bc (${BITCODE_LIB}) is not valid. Is BITCODE_DIR correctly defined?")
set (QUIT 1)
endif()
get_include_path(LLVM_DIR "LLVM directory" RESULT FOUND NAMES "clang" HINTS "${CMAKE_INSTALL_PREFIX}/llvm/bin")
if (NOT ${FOUND})
set (QUIT 1)
endif()
set (CLANG "${LLVM_DIR}/clang")
if (NOT EXISTS ${CLANG})
message("ERROR: path to clang (${CLANG}) is not valid. Is LLVM_DIR correctly defined?")
set (QUIT 1)
endif()
set (LLVM_LINK "${LLVM_DIR}/llvm-link")
if (NOT EXISTS ${LLVM_LINK})
message("ERROR: path to llvm-link (${LLVM_LINK}) is not valid. Is LLVM_DIR correctly defined?")
set (QUIT 1)
endif()
set (LLVM_DIS "${LLVM_DIR}/llvm-dis")
if (NOT EXISTS ${LLVM_DIS})
message("ERROR: path to llvm-dis (${LLVM_DIS}) is not valid. Is LLVM_DIR correctly defined?")
set (QUIT 1)
endif()
# Value of Images Src Dir is bound in parent environment
set (KERNELS_DIR "${IMAGE_SOURCE_DIR}/blit_src")
# Define the target devices with xnack enable
if (NOT DEFINED XNACK_DEVS)
set (XNACK_DEVS "gfx801;gfx902")
@@ -106,29 +75,25 @@ if (QUIT)
return()
endif()
set(TARGET_TRIPLE "amdgcn-amd-amdhsa")
message("")
message("Build Setting:")
message(" Target Devices: ${TARGET_DEVICES}")
message(" Proj. Src Dir: ${PROJECT_SOURCE_DIR}")
message(" Proj. Bld Dir: ${PROJECT_BINARY_DIR}")
message(" Image Source Dir: ${IMAGE_SOURCE_DIR}")
message(" LLVM Dir: ${LLVM_DIR}")
message(" Clang path: ${CLANG}")
message(" OpenCL Dir: ${OPENCL_DIR}")
message(" OpenCL version: ${OPENCL_VER}")
message(" Bitcode Dir: ${BITCODE_DIR}")
message(" Target Triple: ${TARGET_TRIPLE}")
if(${CMAKE_VERBOSE_MAKEFILE})
get_property(clang_path TARGET clang PROPERTY LOCATION)
message("Using clang from: ${clang_path}")
message("Build Setting:")
message(" Target Devices: ${TARGET_DEVICES}")
message(" XNACK Devices: ${XNACK_DEVS}")
message(" Clang path: ${clang_path}")
message(" Bitcode Dir: ${BITCODE_DIR}")
endif()
##==========================================
## Generate Kernel Bitcode
## Add custom command to generate a kernel code object file
##==========================================
function(gen_kernel_bc TARGET_DEV XNACK_OPT FPREFIX INPUT_FILE OUTPUT_FILE)
function(gen_kernel_bc TARGET_DEV XNACK_OPT INPUT_FILE OUTPUT_FILE)
string (REPLACE "gfx" "" GFXIP "${TARGET_DEV}")
separate_arguments(CLANG_ARG_LIST UNIX_COMMAND
"-O2 -x cl -target ${TARGET_TRIPLE} -Xclang -finclude-default-header -mcpu=${TARGET_DEV} -m${XNACK_OPT}
"-O2 -x cl -cl-denorms-are-zero -cl-std=CL2.0 -target amdgcn-amd-amdhsa
-Xclang -finclude-default-header -mcpu=${TARGET_DEV} -m${XNACK_OPT}
-nogpulib
-Xclang -mlink-bitcode-file -Xclang ${BITCODE_DIR}/opencl.amdgcn.bc
-Xclang -mlink-bitcode-file -Xclang ${BITCODE_DIR}/ockl.amdgcn.bc
@@ -137,19 +102,26 @@ function(gen_kernel_bc TARGET_DEV XNACK_OPT FPREFIX INPUT_FILE OUTPUT_FILE)
-Xclang -mlink-bitcode-file -Xclang ${BITCODE_DIR}/oclc_isa_version_${GFXIP}.amdgcn.bc
-Xclang -mlink-bitcode-file -Xclang ${BITCODE_DIR}/oclc_unsafe_math_off.amdgcn.bc
-Xclang -mlink-bitcode-file -Xclang ${BITCODE_DIR}/oclc_finite_only_off.amdgcn.bc
-cl-std=CL${OPENCL_VER} -o ${OUTPUT_FILE} ${INPUT_FILE}")
-o ${OUTPUT_FILE} ${INPUT_FILE}")
add_custom_target("${FPREFIX}" ${CLANG} ${CLANG_ARG_LIST}
COMMENT "BUILDING bitcode for ${FPREFIX}..."
## Add custom command to produce a code object file.
## This depends on the kernel source file & compiler.
## It does not pickup devicelib changes. It is not clear
## how to do that after conversion to --rocm-path is done.
add_custom_command(OUTPUT ${OUTPUT_FILE} COMMAND clang ${CLANG_ARG_LIST}
DEPENDS ${INPUT_FILE} clang
COMMENT "BUILDING bitcode for ${OUTPUT_FILE}..."
VERBATIM)
if(${CMAKE_VERBOSE_MAKEFILE})
message(" Kernel Source: " ${INPUT_FILE})
message(" Kernel Bitcode: " ${OUTPUT_FILE})
endif()
endfunction(gen_kernel_bc)
##==========================================
## Build the kernel for a device
## Find device code object name and forward to custom command
##==========================================
function(build_kernel BLIT_NAME TARG_DEV)
@@ -160,30 +132,33 @@ function(build_kernel BLIT_NAME TARG_DEV)
set (XNACK_OPT "no-xnack")
endif()
set (FILE_PREFIX "${BLIT_NAME}_${TARG_DEV}")
set (HSACO_TARG_LIST ${HSACO_TARG_LIST} "${FILE_PREFIX}" CACHE INTERNAL HSACO_TARG_LIST)
## generate kernel bitcodes
##
set (CL_FILE "${KERNELS_DIR}/imageblit_kernels.cl")
set (KERNEL_BC_FILE "${FILE_PREFIX}")
gen_kernel_bc(${TARG_DEV} ${XNACK_OPT} ${FILE_PREFIX} ${CL_FILE} ${KERNEL_BC_FILE})
set (CODE_OBJECT_FILE "${BLIT_NAME}_${TARG_DEV}")
set (CL_FILE ${CMAKE_CURRENT_SOURCE_DIR}/imageblit_kernels.cl)
gen_kernel_bc(${TARG_DEV} ${XNACK_OPT} ${CL_FILE} ${CODE_OBJECT_FILE})
## Build a list of code object file names
## These will be target dependencies.
set (HSACO_TARG_LIST ${HSACO_TARG_LIST} "${CODE_OBJECT_FILE}" PARENT_SCOPE)
endfunction(build_kernel)
##==========================================
## Build the kernel for a list of devices
##==========================================
function(build_kernel_for_devices BLIT_NAME)
set(HSACO_TARG_LIST PARENT_SCOPE)
set(HSACO_TARG_LIST "")
foreach(dev ${TARGET_DEVICES})
message("\n Working on: ${dev} ...")
if(${CMAKE_VERBOSE_MAKEFILE})
message("\n Generating: ${dev} ...")
endif()
build_kernel(${BLIT_NAME} ${dev})
endforeach(dev)
set(HSACO_TARG_LIST ${HSACO_TARG_LIST} PARENT_SCOPE)
endfunction(build_kernel_for_devices)
##==========================================
@@ -191,13 +166,14 @@ endfunction(build_kernel_for_devices)
##==========================================
function(generate_blit_file BFILE)
file(REMOVE ${IMAGE_SOURCE_DIR}/${BFILE})
## Add a custom command that generates opencl_blit_objects.cpp
## This depends on all the generated code object files and the C++ generator script.
add_custom_command(OUTPUT ${BFILE}
COMMAND ${CMAKE_CURRENT_SOURCE_DIR}/create_hsaco_ascii_file.sh ${CMAKE_CURRENT_BINARY_DIR}/${BFILE}
DEPENDS ${HSACO_TARG_LIST} create_hsaco_ascii_file.sh )
add_custom_command(OUTPUT ${IMAGE_SOURCE_DIR}/${BFILE}
COMMAND ${KERNELS_DIR}/create_hsaco_ascii_file.sh ${IMAGE_SOURCE_DIR}/${BFILE})
message("\n Will create ASCII bitcodes in ${BFILE} for ${TARGET_DEVICES} ... \n")
add_custom_target(${BFILE} DEPENDS ${HSACO_TARG_LIST} ${IMAGE_SOURCE_DIR}/${BFILE})
## Export a target that builds (and depends on) opencl_blit_objects.cpp
add_custom_target( ${BFILE} DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/${BFILE} )
endfunction(generate_blit_file)
projects/rocr-runtime/runtime/hsa-runtime/image/blit_src/create_hsaco_ascii_file.sh
+1 -7
Vedi File
@@ -9,7 +9,6 @@ then
fi
# Create the file in a temporary location and then move it in atomically
rm -rf "$opencl_blit_file.tmp"
{
cat <<EOF
//==============================================================================
@@ -33,9 +32,4 @@ cat <<EOF
EOF
} > "$opencl_blit_file.tmp"
# Move the file atomically into place, so make doesn't get half a file
# but only if it has changed. cmp -s is happy for one file not to exist
cmp -s "$opencl_blit_file.tmp" "$opencl_blit_file" ||
mv -f "$opencl_blit_file.tmp" "$opencl_blit_file"
} > "$opencl_blit_file"