diff --git a/CMakeLists.txt b/CMakeLists.txt index bf42b6aa15..fd321f7722 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -26,11 +26,15 @@ set(version 1.1.0) project(composable_kernel VERSION ${version} LANGUAGES CXX HIP) include(CTest) +# Usage: for customized Python location cmake -DCK_USE_ALTERNATIVE_PYTHON="/opt/Python-3.8.13/bin/python3.8" +# CK Codegen requires dataclass which is added in Python 3.7 +# Python version 3.8 is required for general good practice as it is default for Ubuntu 20.04 if(NOT CK_USE_ALTERNATIVE_PYTHON) - find_package(Python3 3.6 COMPONENTS Interpreter REQUIRED) + find_package(Python3 3.8 COMPONENTS Interpreter REQUIRED) else() message("Using alternative python version") set(EXTRA_PYTHON_PATH) + # this is overly restrictive, we may need to be more flexible on the following string(REPLACE "/bin/python3.8" "" EXTRA_PYTHON_PATH "${CK_USE_ALTERNATIVE_PYTHON}") message("alternative python path is: ${EXTRA_PYTHON_PATH}") find_package(Python3 3.6 COMPONENTS Interpreter REQUIRED) diff --git a/codegen/test/CMakeLists.txt b/codegen/test/CMakeLists.txt index 3ca3f39a30..6dd130bc3f 100644 --- a/codegen/test/CMakeLists.txt +++ b/codegen/test/CMakeLists.txt @@ -6,6 +6,9 @@ if(NOT INSTANCES_ONLY) set_source_files_properties(${TEST_SRC} PROPERTIES LANGUAGE HIP) get_filename_component(BASE_NAME ${TEST_SRC} NAME_WE) add_executable(codegen_test_${BASE_NAME} ${TEST_SRC}) + if(CK_USE_ALTERNATIVE_PYTHON) + target_link_options(codegen_test_${BASE_NAME} PRIVATE -lstdc++fs) + endif() add_dependencies(codegen codegen_test_${BASE_NAME}) add_dependencies(tests codegen_test_${BASE_NAME}) add_dependencies(check codegen_test_${BASE_NAME}) diff --git a/codegen/test/rtc/include/rtc/compile_kernel.hpp b/codegen/test/rtc/include/rtc/compile_kernel.hpp index 5a4a4b0dd6..71db7be249 100644 --- a/codegen/test/rtc/include/rtc/compile_kernel.hpp +++ b/codegen/test/rtc/include/rtc/compile_kernel.hpp @@ -2,14 +2,14 @@ #define GUARD_HOST_TEST_RTC_INCLUDE_RTC_COMPILE_KERNEL #include -#include +#include #include namespace rtc { struct src_file { - std::filesystem::path path; + CK::fs::path path; std::string_view content; }; diff --git a/codegen/test/rtc/include/rtc/tmp_dir.hpp b/codegen/test/rtc/include/rtc/tmp_dir.hpp index f0fd1f72bb..0b4bf002c1 100644 --- a/codegen/test/rtc/include/rtc/tmp_dir.hpp +++ b/codegen/test/rtc/include/rtc/tmp_dir.hpp @@ -2,13 +2,13 @@ #define GUARD_HOST_TEST_RTC_INCLUDE_RTC_TMP_DIR #include -#include +#include namespace rtc { struct tmp_dir { - std::filesystem::path path; + CK::fs::path path; tmp_dir(const std::string& prefix = ""); void execute(const std::string& cmd) const; diff --git a/codegen/test/rtc/src/compile_kernel.cpp b/codegen/test/rtc/src/compile_kernel.cpp index d84ebf4de9..cc1bb80c31 100644 --- a/codegen/test/rtc/src/compile_kernel.cpp +++ b/codegen/test/rtc/src/compile_kernel.cpp @@ -70,9 +70,9 @@ kernel compile_kernel(const std::vector& srcs, compile_options options for(const auto& src : srcs) { - std::filesystem::path full_path = td.path / src.path; - std::filesystem::path parent_path = full_path.parent_path(); - std::filesystem::create_directories(parent_path); + CK::fs::path full_path = td.path / src.path; + CK::fs::path parent_path = full_path.parent_path(); + CK::fs::create_directories(parent_path); write_string(full_path.string(), src.content); if(src.path.extension().string() == ".cpp") { @@ -86,7 +86,7 @@ kernel compile_kernel(const std::vector& srcs, compile_options options td.execute(compiler() + options.flags); auto out_path = td.path / out; - if(not std::filesystem::exists(out_path)) + if(not CK::fs::exists(out_path)) throw std::runtime_error("Output file missing: " + out); auto obj = read_buffer(out_path.string()); diff --git a/codegen/test/rtc/src/tmp_dir.cpp b/codegen/test/rtc/src/tmp_dir.cpp index 1cc8f75b29..659bbbe13f 100644 --- a/codegen/test/rtc/src/tmp_dir.cpp +++ b/codegen/test/rtc/src/tmp_dir.cpp @@ -31,10 +31,10 @@ std::string unique_string(const std::string& prefix) } tmp_dir::tmp_dir(const std::string& prefix) - : path(std::filesystem::temp_directory_path() / + : path(CK::fs::temp_directory_path() / unique_string(prefix.empty() ? "ck-rtc" : "ck-rtc-" + prefix)) { - std::filesystem::create_directories(this->path); + CK::fs::create_directories(this->path); } void tmp_dir::execute(const std::string& cmd) const @@ -43,6 +43,6 @@ void tmp_dir::execute(const std::string& cmd) const std::system(s.c_str()); } -tmp_dir::~tmp_dir() { std::filesystem::remove_all(this->path); } +tmp_dir::~tmp_dir() { CK::fs::remove_all(this->path); } } // namespace rtc diff --git a/example/ck_tile/01_fmha/CMakeLists.txt b/example/ck_tile/01_fmha/CMakeLists.txt index f82207b920..1ba76a523e 100644 --- a/example/ck_tile/01_fmha/CMakeLists.txt +++ b/example/ck_tile/01_fmha/CMakeLists.txt @@ -22,12 +22,20 @@ string(REPLACE ";" "," FMHA_FWD_APIS "${FMHA_FWD_ENABLE_APIS}") execute_process( COMMAND ${Python3_EXECUTABLE} ${CMAKE_CURRENT_LIST_DIR}/generate.py --api ${FMHA_FWD_APIS} --list_blobs ${CMAKE_CURRENT_BINARY_DIR}/fwd_blob_list.txt + RESULT_VARIABLE ret ) +if(ret AND NOT ret EQUAL 0) + message( FATAL_ERROR "CK Tile FMHA FAILED to genrate a list of FWD kernels via Python.") +endif() execute_process( COMMAND ${Python3_EXECUTABLE} ${CMAKE_CURRENT_LIST_DIR}/generate.py --api bwd --list_blobs ${CMAKE_CURRENT_BINARY_DIR}/bwd_blob_list.txt --receipt 3 + RESULT_VARIABLE ret ) +if(ret AND NOT ret EQUAL 0) + message( FATAL_ERROR "CK Tile FMHA FAILED to genrate a list of BWD kernels via Python.") +endif() # NOTE: for cmake, the FMHA_FWD_GEN_BLOBS/FMHA_BWD_GEN_BLOBS files must be in the same directory # as current cmake list, otherwise will not figure out the dependency properly diff --git a/include/ck/filesystem.hpp b/include/ck/filesystem.hpp new file mode 100644 index 0000000000..41ab5f4ddb --- /dev/null +++ b/include/ck/filesystem.hpp @@ -0,0 +1,135 @@ +// SPDX-License-Identifier: MIT +// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved. + +#ifndef GUARD_CK_FILESYSTEM_HPP_ +#define GUARD_CK_FILESYSTEM_HPP_ + +#include +#include + +// clang-format off +#if defined(CPPCHECK) + #define CK_HAS_FILESYSTEM 1 + #define CK_HAS_FILESYSTEM_TS 1 +#elif defined(_WIN32) + #if _MSC_VER >= 1920 + #define CK_HAS_FILESYSTEM 1 + #define CK_HAS_FILESYSTEM_TS 0 + #elif _MSC_VER >= 1900 + #define CK_HAS_FILESYSTEM 0 + #define CK_HAS_FILESYSTEM_TS 1 + #else + #define CK_HAS_FILESYSTEM 0 + #define CK_HAS_FILESYSTEM_TS 0 + #endif +#elif defined(__has_include) + #if __has_include() && __cplusplus >= 201703L + #define CK_HAS_FILESYSTEM 1 + #else + #define CK_HAS_FILESYSTEM 0 + #endif + #if __has_include() && __cplusplus >= 201103L + #define CK_HAS_FILESYSTEM_TS 1 + #else + #define CK_HAS_FILESYSTEM_TS 0 + #endif +#else + #define CK_HAS_FILESYSTEM 0 + #define CK_HAS_FILESYSTEM_TS 0 +#endif +// clang-format on + +#if CK_HAS_FILESYSTEM +#include +#elif CK_HAS_FILESYSTEM_TS +#include +#else +#error "No filesystem include available" +#endif + +namespace CK { + +#if CK_HAS_FILESYSTEM +namespace fs = ::std::filesystem; +#elif CK_HAS_FILESYSTEM_TS +namespace fs = ::std::experimental::filesystem; +#endif + +} // namespace CK + +inline std::string operator+(const std::string_view s, const CK::fs::path& path) +{ + return path.string().insert(0, s); +} + +inline std::string operator+(const CK::fs::path& path, const std::string_view s) +{ + return path.string().append(s); +} + +#define FS_ENUM_PERMS_ALL fs::perms::all + +#if CK_HAS_FILESYSTEM_TS +#ifdef __linux__ +#include +namespace CK { +inline fs::path weakly_canonical(const fs::path& path) +{ + std::string result(PATH_MAX, '\0'); + std::string p{path.is_relative() ? (fs::current_path() / path).string() : path.string()}; + char* retval = realpath(p.c_str(), &result[0]); + return (retval == nullptr) ? path : fs::path{result}; +} +} // namespace CK +#else +#error "Not implmeneted!" +#endif +#else +namespace CK { +inline fs::path weakly_canonical(const fs::path& path) { return fs::weakly_canonical(path); } +} // namespace CK +#endif + +namespace CK { + +#ifdef _WIN32 +constexpr std::string_view executable_postfix{".exe"}; +constexpr std::string_view library_prefix{""}; +constexpr std::string_view dynamic_library_postfix{".dll"}; +constexpr std::string_view static_library_postfix{".lib"}; +constexpr std::string_view object_file_postfix{".obj"}; +#else +constexpr std::string_view executable_postfix{""}; +constexpr std::string_view library_prefix{"lib"}; +constexpr std::string_view dynamic_library_postfix{".so"}; +constexpr std::string_view static_library_postfix{".a"}; +constexpr std::string_view object_file_postfix{".o"}; +#endif + +inline fs::path make_executable_name(const fs::path& path) +{ + return path.parent_path() / (path.filename() + executable_postfix); +} + +inline fs::path make_dynamic_library_name(const fs::path& path) +{ + return path.parent_path() / (library_prefix + path.filename() + dynamic_library_postfix); +} + +inline fs::path make_object_file_name(const fs::path& path) +{ + return path.parent_path() / (path.filename() + object_file_postfix); +} + +inline fs::path make_static_library_name(const fs::path& path) +{ + return path.parent_path() / (library_prefix + path.filename() + static_library_postfix); +} + +struct FsPathHash +{ + std::size_t operator()(const fs::path& path) const { return fs::hash_value(path); } +}; +} // namespace CK + +#endif // GUARD_CK_FILESYSTEM_HPP_ diff --git a/library/src/tensor_operation_instance/gpu/mha/CMakeLists.txt b/library/src/tensor_operation_instance/gpu/mha/CMakeLists.txt index 3499779c2a..8bfda97e71 100644 --- a/library/src/tensor_operation_instance/gpu/mha/CMakeLists.txt +++ b/library/src/tensor_operation_instance/gpu/mha/CMakeLists.txt @@ -1,19 +1,23 @@ set(FMHA_CPP_FOLDER ${CMAKE_CURRENT_BINARY_DIR}) set(FMHA_SRC_FOLDER ${CMAKE_SOURCE_DIR}/example/ck_tile/01_fmha/) set(CK_TILE_SRC_FOLDER ${CMAKE_SOURCE_DIR}/include/ck_tile/) -# python stuff + +# Usage: for customized Python location cmake -DCK_USE_ALTERNATIVE_PYTHON="/opt/Python-3.8.13/bin/python3.8" +# CK Codegen requires dataclass which is added in Python 3.7 +# Python version 3.8 is required for general good practice as it is default for Ubuntu 20.04 if(NOT CK_USE_ALTERNATIVE_PYTHON) - find_package(PythonInterp 3 REQUIRED) + find_package(PythonInterp 3 REQUIRED) else() - message("Using alternative python version") - set(EXTRA_PYTHON_PATH) - string(REPLACE "/bin/python3.8" "" EXTRA_PYTHON_PATH "${CK_USE_ALTERNATIVE_PYTHON}") - message("alternative python path is: ${EXTRA_PYTHON_PATH}") - find_package(Python3 3.6 COMPONENTS Interpreter REQUIRED) - add_definitions(-DPython3_EXECUTABLE="${CK_USE_ALTERNATIVE_PYTHON}") - set(Python3_EXECUTABLE "${CK_USE_ALTERNATIVE_PYTHON}") - set(PYTHON_EXECUTABLE "${CK_USE_ALTERNATIVE_PYTHON}") - set(ENV{LD_LIBRARY_PATH} "${EXTRA_PYTHON_PATH}/lib:$ENV{LD_LIBRARY_PATH}") + message("Using alternative python version") + set(EXTRA_PYTHON_PATH) + # this is overly restrictive, we may need to be more flexible on the following + string(REPLACE "/bin/python3.8" "" EXTRA_PYTHON_PATH "${CK_USE_ALTERNATIVE_PYTHON}") + message("alternative python path is: ${EXTRA_PYTHON_PATH}") + find_package(Python3 3.6 COMPONENTS Interpreter REQUIRED) + add_definitions(-DPython3_EXECUTABLE="${CK_USE_ALTERNATIVE_PYTHON}") + set(Python3_EXECUTABLE "${CK_USE_ALTERNATIVE_PYTHON}") + set(PYTHON_EXECUTABLE "${CK_USE_ALTERNATIVE_PYTHON}") + set(ENV{LD_LIBRARY_PATH} "${EXTRA_PYTHON_PATH}/lib:$ENV{LD_LIBRARY_PATH}") endif() rocm_install(DIRECTORY ${CK_TILE_SRC_FOLDER} DESTINATION ${CMAKE_INSTALL_INCLUDEDIR}/ck_tile) @@ -27,8 +31,13 @@ file(COPY ${MHA_HEADERS} DESTINATION ${FMHA_CPP_FOLDER}) execute_process( COMMAND ${PYTHON_EXECUTABLE} ${CMAKE_SOURCE_DIR}/example/ck_tile/01_fmha/generate.py --list_blobs ${FMHA_CPP_FOLDER}/blob_list.txt + RESULT_VARIABLE ret ) -file(STRINGS ${FMHA_CPP_FOLDER}/blob_list.txt FMHA_FWD_GEN_BLOBS) +if(ret AND NOT ret EQUAL 0) + message( FATAL_ERROR "CK Tile MHA FAILED to genrate a list of kernels via Python.") +else() + file(STRINGS ${FMHA_CPP_FOLDER}/blob_list.txt FMHA_FWD_GEN_BLOBS) +endif() # actually generate the cpp files add_custom_command( @@ -52,8 +61,6 @@ add_custom_target(generate_cpp_files DEPENDS ${FMHA_FWD_GEN_BLOBS}) add_instance_library(device_mha_instance ${device_files}) - - if (TARGET device_mha_instance) add_dependencies(device_mha_instance generate_cpp_files) endif()