From 7ffe5cdd532ad3abde2f2695f3d270762599a69d Mon Sep 17 00:00:00 2001 From: Jing Zhou Date: Mon, 27 Apr 2020 10:05:22 -0700 Subject: [PATCH 1/6] update tensile version --- CMakeLists.txt | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index ccb7e5b..014a409 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -48,7 +48,7 @@ if(BUILD_DEV) set(MIOPEN_TENSILE_SRC dev) endif() # Use the virtual-env setup and download package from specified repo: -set( MIOPEN_TENSILE_TAG fdd9ef8d5a0687596efee85b7ec187f1fb097087 CACHE STRING "Tensile tag to download" ) +set( MIOPEN_TENSILE_TAG 698e3174c4d6e5045e9a619158cd075273ccef96 CACHE STRING "Tensile tag to download" ) virtualenv_install(wheel) virtualenv_install("git+https://github.com/ROCmSoftwarePlatform/Tensile.git@${MIOPEN_TENSILE_TAG}") list(APPEND CMAKE_PREFIX_PATH ${VIRTUALENV_HOME_DIR}) From c097a321b1a428c3a620df2fd97a3e3b6d3bc938 Mon Sep 17 00:00:00 2001 From: Jing Zhou Date: Wed, 29 Apr 2020 17:52:37 -0700 Subject: [PATCH 2/6] packaging --- CMakeLists.txt | 21 ++++++- install.sh | 158 +++++++++++++++++++++++++++++++++++++++++++++++++ 2 files changed, 178 insertions(+), 1 deletion(-) create mode 100755 install.sh diff --git a/CMakeLists.txt b/CMakeLists.txt index 014a409..3f4a4a2 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -13,6 +13,7 @@ include(ROCMSetupVersion) rocm_setup_version(VERSION 1.0) +set(CMAKE_INSTALL_PREFIX "/usr/local" CACHE PATH "") set(CMAKE_LIBRARY_OUTPUT_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR}/lib) set(CMAKE_ARCHIVE_OUTPUT_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR}/lib) set(CMAKE_RUNTIME_OUTPUT_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR}/bin) @@ -25,6 +26,8 @@ else() set(Tensile_ARCHITECTURE ${AMDGPU_TARGETS}) endif() +# Build options +option(BUILD_TEST "Build tests" OFF) # Dont build as shared as a workaround # option( BUILD_SHARED_LIBS "Build as a shared library" ON ) @@ -95,5 +98,21 @@ rocm_export_targets( install(DIRECTORY "${CMAKE_CURRENT_BINARY_DIR}/lib/miopentensile" DESTINATION lib) -add_subdirectory(test) +# Tests +if(BUILD_TEST) + enable_testing() + add_subdirectory(test) +endif() + +# Package +set(CPACK_DEBIAN_PACKAGE_DEPENDS "rocm-dev (>= 2.5.27)") +set(CPACK_RPM_PACKAGE_REQUIRES "rocm-dev >= 2.5.27") +if(NOT CPACK_PACKAGING_INSTALL_PREFIX) + set(CPACK_PACKAGING_INSTALL_PREFIX "${CMAKE_INSTALL_PREFIX}") +endif() +rocm_create_package( + NAME MIOpenTensile + DESCRIPTION "MIOpen-Tensile Tool" + MAINTAINER "Paul Fultz II " +) diff --git a/install.sh b/install.sh new file mode 100755 index 0000000..9de1274 --- /dev/null +++ b/install.sh @@ -0,0 +1,158 @@ +#!/bin/bash + + +# ################################################# +# helper functions +# ################################################# +function display_help() +{ + echo "rocThust build & installation helper script" + echo "./install [-h|--help] " + echo " [-h|--help] prints this help message" + echo " [-i|--install] install after build" + echo " [-p|--package] build package" + echo " [-r]--relocatable] create a package to support relocatable ROCm" + #Not implemented yet + # echo " [-d|--dependencies] install build dependencies" + echo " [-c|--clients] build library clients too (combines with -i & -d)" + echo " [-g|--debug] -DCMAKE_BUILD_TYPE=Debug (default is =Release)" + echo " [--hip-clang] build library for amdgpu backend using hip-clang" +} + + +# ################################################# +# global variables +# ################################################# +install_package=false +build_package=false +build_clients=false +build_release=true +build_type=Release +build_hip_clang=false +run_tests=false +rocm_path=/opt/rocm +build_relocatable=false + +# ################################################# +# Parameter parsing +# ################################################# + +# check if we have a modern version of getopt that can handle whitespace and long parameters +getopt -T +if [[ $? -eq 4 ]]; then + GETOPT_PARSE=$(getopt --name "${0}" --longoptions help,install,clients,debug,hip-clang,test,package,relocatable --options hicdtprg -- "$@") +else + echo "Need a new version of getopt" + exit 1 +fi + +if [[ $? -ne 0 ]]; then + echo "getopt invocation failed; could not parse the command line"; + exit 1 +fi + +eval set -- "${GETOPT_PARSE}" + +check_exit_code( ) +{ + if (( $1 != 0 )); then + exit $1 + fi +} + +while true; do + case "${1}" in + -h|--help) + display_help + exit 0 + ;; + -i|--install) + install_package=true + shift ;; + -p|--package) + build_package=true + shift ;; + -c|--clients) + build_clients=true + shift ;; + -r|--relocatable) + build_relocatable=true + shift ;; + -g|--debug) + build_type=Debug + build_release=false + shift ;; + -t|--test) + run_tests=true + shift ;; + --hip-clang) + build_hip_clang=true + shift ;; + --) shift ; break ;; + *) echo "Unexpected command line parameter received; aborting"; + exit 1 + ;; + esac +done + +if [[ "${build_relocatable}" == true ]]; then + if ! [ -z ${ROCM_PATH+x} ]; then + rocm_path=${ROCM_PATH} + fi +fi + +# Create and go to the build directory. +mkdir -p build; cd build + +if ($build_release); then + mkdir -p release; cd release +else + mkdir -p debug; cd debug +fi + +# Set compiler +compiler="hcc" +if [[ "${build_hip_clang}" == true ]]; then + compiler="hipcc" +fi + +cmake_executable="cmake" +if [ -e /etc/redhat-release ] ; then + cmake_executable="cmake3" +fi + +build_test="OFF" +if [[ "${build_clients}" == true ]]; then + build_test="ON" +fi + +if [[ "${build_relocatable}" == true ]]; then + CXX=$rocm_path/bin/${compiler} ${cmake_executable} \ + -DCMAKE_INSTALL_PREFIX=${rocm_path} \ + -DCMAKE_PREFIX_PATH="${rocm_path} ${rocm_path}/hcc ${rocm_path}/hip" \ + -DCMAKE_MODULE_PATH="${rocm_path}/hip/cmake" \ + -DBUILD_TEST=${build_test} \ + ../../. # or cmake-gui ../. +else + CXX=$rocm_path/bin/${compiler} ${cmake_executable} -DBUILD_TEST=${build_test} ../../. # or cmake-gui ../. +fi +check_exit_code "$?" + +# Build +make -j$(nproc) +check_exit_code "$?" + +if ($run_tests); then +# Optionally, run tests if they're enabled. +ctest --output-on-failure +fi + +if ($install_package); then + make install + check_exit_code "$?" +fi + +if ($build_package); then + make package -j$(nproc) + check_exit_code "$?" +fi From 6a57d799b42ba46f398bb2dd078f9c161c8ae219 Mon Sep 17 00:00:00 2001 From: Paul Date: Wed, 29 Apr 2020 23:22:31 -0500 Subject: [PATCH 3/6] Remove install script --- CMakeLists.txt | 25 +------- install.sh | 158 ------------------------------------------------- 2 files changed, 2 insertions(+), 181 deletions(-) delete mode 100755 install.sh diff --git a/CMakeLists.txt b/CMakeLists.txt index 3f4a4a2..b82a43b 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -13,7 +13,7 @@ include(ROCMSetupVersion) rocm_setup_version(VERSION 1.0) -set(CMAKE_INSTALL_PREFIX "/usr/local" CACHE PATH "") +set(CMAKE_INSTALL_PREFIX "/opt/rocm" CACHE PATH "") set(CMAKE_LIBRARY_OUTPUT_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR}/lib) set(CMAKE_ARCHIVE_OUTPUT_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR}/lib) set(CMAKE_RUNTIME_OUTPUT_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR}/bin) @@ -26,8 +26,6 @@ else() set(Tensile_ARCHITECTURE ${AMDGPU_TARGETS}) endif() -# Build options -option(BUILD_TEST "Build tests" OFF) # Dont build as shared as a workaround # option( BUILD_SHARED_LIBS "Build as a shared library" ON ) @@ -96,23 +94,4 @@ rocm_export_targets( hip ) -install(DIRECTORY "${CMAKE_CURRENT_BINARY_DIR}/lib/miopentensile" DESTINATION lib) - -# Tests -if(BUILD_TEST) - enable_testing() - add_subdirectory(test) -endif() - -# Package -set(CPACK_DEBIAN_PACKAGE_DEPENDS "rocm-dev (>= 2.5.27)") -set(CPACK_RPM_PACKAGE_REQUIRES "rocm-dev >= 2.5.27") -if(NOT CPACK_PACKAGING_INSTALL_PREFIX) - set(CPACK_PACKAGING_INSTALL_PREFIX "${CMAKE_INSTALL_PREFIX}") -endif() - -rocm_create_package( - NAME MIOpenTensile - DESCRIPTION "MIOpen-Tensile Tool" - MAINTAINER "Paul Fultz II " -) +add_subdirectory(test) diff --git a/install.sh b/install.sh deleted file mode 100755 index 9de1274..0000000 --- a/install.sh +++ /dev/null @@ -1,158 +0,0 @@ -#!/bin/bash - - -# ################################################# -# helper functions -# ################################################# -function display_help() -{ - echo "rocThust build & installation helper script" - echo "./install [-h|--help] " - echo " [-h|--help] prints this help message" - echo " [-i|--install] install after build" - echo " [-p|--package] build package" - echo " [-r]--relocatable] create a package to support relocatable ROCm" - #Not implemented yet - # echo " [-d|--dependencies] install build dependencies" - echo " [-c|--clients] build library clients too (combines with -i & -d)" - echo " [-g|--debug] -DCMAKE_BUILD_TYPE=Debug (default is =Release)" - echo " [--hip-clang] build library for amdgpu backend using hip-clang" -} - - -# ################################################# -# global variables -# ################################################# -install_package=false -build_package=false -build_clients=false -build_release=true -build_type=Release -build_hip_clang=false -run_tests=false -rocm_path=/opt/rocm -build_relocatable=false - -# ################################################# -# Parameter parsing -# ################################################# - -# check if we have a modern version of getopt that can handle whitespace and long parameters -getopt -T -if [[ $? -eq 4 ]]; then - GETOPT_PARSE=$(getopt --name "${0}" --longoptions help,install,clients,debug,hip-clang,test,package,relocatable --options hicdtprg -- "$@") -else - echo "Need a new version of getopt" - exit 1 -fi - -if [[ $? -ne 0 ]]; then - echo "getopt invocation failed; could not parse the command line"; - exit 1 -fi - -eval set -- "${GETOPT_PARSE}" - -check_exit_code( ) -{ - if (( $1 != 0 )); then - exit $1 - fi -} - -while true; do - case "${1}" in - -h|--help) - display_help - exit 0 - ;; - -i|--install) - install_package=true - shift ;; - -p|--package) - build_package=true - shift ;; - -c|--clients) - build_clients=true - shift ;; - -r|--relocatable) - build_relocatable=true - shift ;; - -g|--debug) - build_type=Debug - build_release=false - shift ;; - -t|--test) - run_tests=true - shift ;; - --hip-clang) - build_hip_clang=true - shift ;; - --) shift ; break ;; - *) echo "Unexpected command line parameter received; aborting"; - exit 1 - ;; - esac -done - -if [[ "${build_relocatable}" == true ]]; then - if ! [ -z ${ROCM_PATH+x} ]; then - rocm_path=${ROCM_PATH} - fi -fi - -# Create and go to the build directory. -mkdir -p build; cd build - -if ($build_release); then - mkdir -p release; cd release -else - mkdir -p debug; cd debug -fi - -# Set compiler -compiler="hcc" -if [[ "${build_hip_clang}" == true ]]; then - compiler="hipcc" -fi - -cmake_executable="cmake" -if [ -e /etc/redhat-release ] ; then - cmake_executable="cmake3" -fi - -build_test="OFF" -if [[ "${build_clients}" == true ]]; then - build_test="ON" -fi - -if [[ "${build_relocatable}" == true ]]; then - CXX=$rocm_path/bin/${compiler} ${cmake_executable} \ - -DCMAKE_INSTALL_PREFIX=${rocm_path} \ - -DCMAKE_PREFIX_PATH="${rocm_path} ${rocm_path}/hcc ${rocm_path}/hip" \ - -DCMAKE_MODULE_PATH="${rocm_path}/hip/cmake" \ - -DBUILD_TEST=${build_test} \ - ../../. # or cmake-gui ../. -else - CXX=$rocm_path/bin/${compiler} ${cmake_executable} -DBUILD_TEST=${build_test} ../../. # or cmake-gui ../. -fi -check_exit_code "$?" - -# Build -make -j$(nproc) -check_exit_code "$?" - -if ($run_tests); then -# Optionally, run tests if they're enabled. -ctest --output-on-failure -fi - -if ($install_package); then - make install - check_exit_code "$?" -fi - -if ($build_package); then - make package -j$(nproc) - check_exit_code "$?" -fi From c0f0d67a361179a7084720abfd1ee9f5f5b38665 Mon Sep 17 00:00:00 2001 From: Paul Date: Thu, 30 Apr 2020 15:43:35 -0500 Subject: [PATCH 4/6] Add back the installation for the code objects --- CMakeLists.txt | 2 ++ 1 file changed, 2 insertions(+) diff --git a/CMakeLists.txt b/CMakeLists.txt index b82a43b..bab3308 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -73,6 +73,8 @@ target_link_libraries(MIOpenTensile PUBLIC hip::host) target_link_libraries(MIOpenTensile PRIVATE TensileHost) target_compile_definitions(MIOpenTensile PRIVATE __HIP_PLATFORM_HCC__) +install(DIRECTORY "${CMAKE_CURRENT_BINARY_DIR}/lib/miopentensile" DESTINATION lib) + include(ROCMCreatePackage) rocm_create_package( NAME MIOpenTensile From 1d02871713020eafad1992cb5cf52afe08a11fe2 Mon Sep 17 00:00:00 2001 From: Paul Date: Thu, 30 Apr 2020 20:48:05 -0500 Subject: [PATCH 5/6] Set install prefix --- CMakeLists.txt | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index bab3308..ad7a088 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -1,6 +1,8 @@ cmake_minimum_required(VERSION 3.5) +set(CMAKE_INSTALL_PREFIX "/opt/rocm" CACHE PATH "") + project(miopentensile LANGUAGES CXX) add_definitions(-std=c++14) @@ -13,7 +15,6 @@ include(ROCMSetupVersion) rocm_setup_version(VERSION 1.0) -set(CMAKE_INSTALL_PREFIX "/opt/rocm" CACHE PATH "") set(CMAKE_LIBRARY_OUTPUT_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR}/lib) set(CMAKE_ARCHIVE_OUTPUT_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR}/lib) set(CMAKE_RUNTIME_OUTPUT_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR}/bin) From 32396de80fcd92ed55a829e78be2367c62212285 Mon Sep 17 00:00:00 2001 From: Paul Date: Wed, 6 May 2020 11:24:58 -0500 Subject: [PATCH 6/6] Embed code objects --- CMakeLists.txt | 8 +++- cmake/embed.cmake | 97 +++++++++++++++++++++++++++++++++++++++++++++++ src/gemm_api.cpp | 5 ++- 3 files changed, 106 insertions(+), 4 deletions(-) create mode 100644 cmake/embed.cmake diff --git a/CMakeLists.txt b/CMakeLists.txt index ad7a088..f4b8461 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -68,13 +68,17 @@ TensileCreateLibraryFiles( NO_MERGE_FILES ) +include(embed) + +file(GLOB CODE_OBJECTS "${CMAKE_CURRENT_BINARY_DIR}/lib/miopentensile/library/*co") +add_embed_library(miopentensile_data ${CODE_OBJECTS}) add_library(MIOpenTensile SHARED src/gemm_api.cpp) target_link_libraries(MIOpenTensile PUBLIC hip::host) -target_link_libraries(MIOpenTensile PRIVATE TensileHost) +target_link_libraries(MIOpenTensile PRIVATE TensileHost miopentensile_data) target_compile_definitions(MIOpenTensile PRIVATE __HIP_PLATFORM_HCC__) -install(DIRECTORY "${CMAKE_CURRENT_BINARY_DIR}/lib/miopentensile" DESTINATION lib) +install(FILES "${CMAKE_CURRENT_BINARY_DIR}/lib/miopentensile/library/TensileLibrary.yaml" DESTINATION lib/miopentensile/library) include(ROCMCreatePackage) rocm_create_package( diff --git a/cmake/embed.cmake b/cmake/embed.cmake new file mode 100644 index 0000000..b388a4d --- /dev/null +++ b/cmake/embed.cmake @@ -0,0 +1,97 @@ + +find_program(EMBED_LD ld) +find_program(EMBED_OBJCOPY objcopy) + +function(generate_embed_source EMBED_NAME) + set(options) + set(oneValueArgs SRC HEADER) + set(multiValueArgs OBJECTS SYMBOLS) + + cmake_parse_arguments(PARSE "${options}" "${oneValueArgs}" "${multiValueArgs}" ${ARGN}) + + set(EXTERNS) + set(INIT_KERNELS) + + # TODO: check the lens are the same + list(LENGTH PARSE_SYMBOLS SYMBOLS_LEN) + list(LENGTH PARSE_OBJECTS OBJECTS_LEN) + if(NOT ${SYMBOLS_LEN} EQUAL ${OBJECTS_LEN}) + message(FATAL_ERROR "Symbols and objects dont match: ${SYMBOLS_LEN} != ${OBJECTS_LEN}") + endif() + math(EXPR LEN "${SYMBOLS_LEN} - 1") + + foreach(idx RANGE ${LEN}) + list(GET PARSE_SYMBOLS ${idx} SYMBOL) + list(GET PARSE_OBJECTS ${idx} OBJECT) + # foreach(SYMBOL ${PARSE_SYMBOLS}) + set(START_SYMBOL "_binary_${SYMBOL}_start") + set(END_SYMBOL "_binary_${SYMBOL}_end") + string(APPEND EXTERNS " + extern const char ${START_SYMBOL}[]; + extern const char ${END_SYMBOL}[]; + ") + + get_filename_component(BASE_NAME "${OBJECT}" NAME) + + string(APPEND INIT_KERNELS " + { \"${BASE_NAME}\", { ${START_SYMBOL}, ${END_SYMBOL}} }, + ") + endforeach() + + file(WRITE "${PARSE_HEADER}" " +#include +const std::unordered_map>& ${EMBED_NAME}(); +") + + file(WRITE "${PARSE_SRC}" " +#include <${EMBED_NAME}.hpp> +${EXTERNS} +const std::unordered_map>& ${EMBED_NAME}() +{ + static const std::unordered_map> result = {${INIT_KERNELS}}; + return result; +} +") +endfunction() + +function(embed_file OUTPUT_FILE OUTPUT_SYMBOL FILE) + set(${OUTPUT_FILE} "${FILE}.o" PARENT_SCOPE) + set(WORKING_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR}) + # Glob is used to compute the relative path + file(GLOB FILES RELATIVE ${WORKING_DIRECTORY} ${FILE}) + foreach(REL_FILE ${FILES}) + string(MAKE_C_IDENTIFIER "${REL_FILE}" SYMBOL) + set(${OUTPUT_SYMBOL} ${SYMBOL} PARENT_SCOPE) + add_custom_command( + OUTPUT "${FILE}.o" + COMMAND ${EMBED_LD} -r -o "${REL_FILE}.o" -z noexecstack --format=binary "${REL_FILE}" + COMMAND ${EMBED_OBJCOPY} --rename-section .data=.rodata,alloc,load,readonly,data,contents "${REL_FILE}.o" + WORKING_DIRECTORY ${WORKING_DIRECTORY} + DEPENDS ${FILE} + VERBATIM + ) + endforeach() +endfunction() + +function(add_embed_library EMBED_NAME) + file(MAKE_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR}/embed) + file(MAKE_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR}/embed/${EMBED_NAME}) + set(EMBED_DIR ${CMAKE_CURRENT_BINARY_DIR}/embed/${EMBED_NAME}) + set(SRC_FILE "${EMBED_DIR}/${EMBED_NAME}.cpp") + set(HEADER_FILE "${EMBED_DIR}/include/${EMBED_NAME}.hpp") + set(WORKING_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR}) + set(OUTPUT_FILES) + set(SYMBOLS) + message(STATUS "Embedding files") + foreach(FILE ${ARGN}) + # message(STATUS "Embedding file ${FILE}") + embed_file(OUTPUT_FILE OUTPUT_SYMBOL ${FILE}) + list(APPEND OUTPUT_FILES ${OUTPUT_FILE}) + list(APPEND SYMBOLS ${OUTPUT_SYMBOL}) + endforeach() + message(STATUS "Generating embedding library ${EMBED_NAME}") + generate_embed_source(${EMBED_NAME} SRC ${SRC_FILE} HEADER ${HEADER_FILE} OBJECTS ${OUTPUT_FILES} SYMBOLS ${SYMBOLS}) + add_library(${EMBED_NAME} STATIC ${OUTPUT_FILES} "${SRC_FILE}") + target_include_directories(${EMBED_NAME} PUBLIC "${EMBED_DIR}/include") + set_target_properties(${EMBED_NAME} PROPERTIES POSITION_INDEPENDENT_CODE On) +endfunction() diff --git a/src/gemm_api.cpp b/src/gemm_api.cpp index a86fd79..33defee 100644 --- a/src/gemm_api.cpp +++ b/src/gemm_api.cpp @@ -6,6 +6,7 @@ #include #include #include +#include std::vector glob_files(const std::string& s) { @@ -64,8 +65,8 @@ const auto& library() auto create_adaptor() { // Workaround: The Tensile::hip::SolutionAdapter is not a regular type, so heap allocate it instead auto a = std::make_shared(); - for(auto&& f:glob_files(library_path() + "*co")) - a->loadCodeObjectFile(f); + for(auto&& p:miopentensile_data()) + a->loadCodeObject(p.second.first); return a; }