diff --git a/projects/rocr-runtime/runtime/hsa-runtime/CMakeLists.txt b/projects/rocr-runtime/runtime/hsa-runtime/CMakeLists.txt index f9120455c6..ba52277b9e 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/CMakeLists.txt +++ b/projects/rocr-runtime/runtime/hsa-runtime/CMakeLists.txt @@ -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 + $ + $ + 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} $ ) -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" ) diff --git a/projects/rocr-runtime/runtime/hsa-runtime/DEBIAN/post_install b/projects/rocr-runtime/runtime/hsa-runtime/DEBIAN/post_install index a9afa9c26c..e1f3fe3ac2 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/DEBIAN/post_install +++ b/projects/rocr-runtime/runtime/hsa-runtime/DEBIAN/post_install @@ -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" diff --git a/projects/rocr-runtime/runtime/hsa-runtime/DEBIAN/pre_remove b/projects/rocr-runtime/runtime/hsa-runtime/DEBIAN/pre_remove index 2dd27fdc36..3911851494 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/DEBIAN/pre_remove +++ b/projects/rocr-runtime/runtime/hsa-runtime/DEBIAN/pre_remove @@ -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) ;; diff --git a/projects/rocr-runtime/runtime/hsa-runtime/cmake_modules/hsa_common.cmake b/projects/rocr-runtime/runtime/hsa-runtime/cmake_modules/hsa_common.cmake index 3bf1c1a9ae..7e8b267346 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/cmake_modules/hsa_common.cmake +++ b/projects/rocr-runtime/runtime/hsa-runtime/cmake_modules/hsa_common.cmake @@ -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") diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/hsa_ext_interface.cpp b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/hsa_ext_interface.cpp index b9b814be09..b6a8dbeb02 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/hsa_ext_interface.cpp +++ b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/hsa_ext_interface.cpp @@ -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; } diff --git a/projects/rocr-runtime/runtime/hsa-runtime/hsa-runtime64-config.cmake.in b/projects/rocr-runtime/runtime/hsa-runtime/hsa-runtime64-config.cmake.in new file mode 100644 index 0000000000..41da4f949c --- /dev/null +++ b/projects/rocr-runtime/runtime/hsa-runtime/hsa-runtime64-config.cmake.in @@ -0,0 +1,7 @@ +@PACKAGE_INIT@ + +include( CMakeFindDependencyMacro ) + +find_dependency(hsakmt 1.0) + +include( "${CMAKE_CURRENT_LIST_DIR}/@CORE_RUNTIME_TARGET@Targets.cmake" ) diff --git a/projects/rocr-runtime/runtime/hsa-runtime/image/CMakeLists.txt b/projects/rocr-runtime/runtime/hsa-runtime/image/CMakeLists.txt deleted file mode 100755 index c7f31cff97..0000000000 --- a/projects/rocr-runtime/runtime/hsa-runtime/image/CMakeLists.txt +++ /dev/null @@ -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} $ ) -#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() - - diff --git a/projects/rocr-runtime/runtime/hsa-runtime/image/blit_kernel_cl.cl b/projects/rocr-runtime/runtime/hsa-runtime/image/blit_kernel_cl.cl deleted file mode 100644 index 12f5a86abd..0000000000 --- a/projects/rocr-runtime/runtime/hsa-runtime/image/blit_kernel_cl.cl +++ /dev/null @@ -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; - } -} \ No newline at end of file diff --git a/projects/rocr-runtime/runtime/hsa-runtime/image/blit_src/CMakeLists.txt b/projects/rocr-runtime/runtime/hsa-runtime/image/blit_src/CMakeLists.txt index f2e025f172..5380325846 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/image/blit_src/CMakeLists.txt +++ b/projects/rocr-runtime/runtime/hsa-runtime/image/blit_src/CMakeLists.txt @@ -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 "/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) diff --git a/projects/rocr-runtime/runtime/hsa-runtime/image/blit_src/create_hsaco_ascii_file.sh b/projects/rocr-runtime/runtime/hsa-runtime/image/blit_src/create_hsaco_ascii_file.sh index 647fbb492d..0f1544d1dc 100755 --- a/projects/rocr-runtime/runtime/hsa-runtime/image/blit_src/create_hsaco_ascii_file.sh +++ b/projects/rocr-runtime/runtime/hsa-runtime/image/blit_src/create_hsaco_ascii_file.sh @@ -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 < "$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"