CMakeLists.txt

set(CMAKE_CXX_COMPILER /opt/rocm/llvm/bin/clang++)

## for online-compiling of HIP kernels
set(OLC_HIP_COMPILER ${CMAKE_CXX_COMPILER} CACHE PATH "")

## reset to avoid the C++ options from the parent project
set(CMAKE_CXX_FLAGS "")
message("Compiling options for library and kernels: ${CMAKE_CXX_FLAGS}")

# look for and register clang-offload-bundler
if(OLC_HIP_COMPILER MATCHES ".*clang\\+\\+$")
    find_program(OLC_OFFLOADBUNDLER_BIN clang-offload-bundler
        PATH_SUFFIXES bin
        PATHS
	    /opt/rocm/llvm
	    ${CMAKE_INSTALL_PREFIX}/llvm
    )
endif()

if(OLC_OFFLOADBUNDLER_BIN)
    message(STATUS "clang-offload-bundler found: ${OLC_OFFLOADBUNDLER_BIN}")
    set(OLC_OFFLOADBUNDLER_BIN "${OLC_OFFLOADBUNDLER_BIN}")
else()
    # look for and register extractkernel
    message(STATUS "clang-offload-bundler not found")

    find_program(EXTRACTKERNEL_BIN extractkernel
        PATH_SUFFIXES bin
        PATHS
            /opt/rocm/hip
            /opt/rocm/hcc
            /opt/rocm
	    ${CMAKE_INSTALL_PREFIX}/hip
            ${CMAKE_INSTALL_PREFIX}/hcc
            ${CMAKE_INSTALL_PREFIX}

    )
    if(EXTRACTKERNEL_BIN)
        message(STATUS "extractkernel found: ${EXTRACTKERNEL_BIN}")
        set(EXTRACTKERNEL_BIN "${EXTRACTKERNEL_BIN}")
    else()
        message(FATAL_ERROR "extractkernel not found")
    endif()
endif()

option(Boost_USE_STATIC_LIBS "Use boost static libraries" OFF)
set(BOOST_COMPONENTS filesystem)
add_definitions(-DBOOST_ALL_NO_LIB=1)
find_package(Boost REQUIRED COMPONENTS ${BOOST_COMPONENTS})

# HIP is always required
find_package(hip REQUIRED PATHS /opt/rocm)
message(STATUS "Build with HIP ${hip_VERSION}")
target_flags(HIP_COMPILER_FLAGS hip::device)
# Remove cuda arch flags
string(REGEX REPLACE --cuda-gpu-arch=[a-z0-9]+ "" HIP_COMPILER_FLAGS "${HIP_COMPILER_FLAGS}")
string(REGEX REPLACE --offload-arch=[a-z0-9]+ "" HIP_COMPILER_FLAGS "${HIP_COMPILER_FLAGS}")

set(OLC_hip_VERSION_MAJOR "${hip_VERSION_MAJOR}")
set(OLC_hip_VERSION_MINOR "${hip_VERSION_MINOR}")
set(OLC_hip_VERSION_PATCH "${hip_VERSION_PATCH}")

option(ENABLE_DEBUG "Build to enable debugging" ON)
if(ENABLE_DEBUG)
    set(OLC_DEBUG 1)
else()
    set(OLC_DEBUG 0)
endif()

configure_file("${PROJECT_SOURCE_DIR}/host/online_compile/include/config.h.in" "${PROJECT_BINARY_DIR}/host/online_compile/include/config.h")

include_directories(BEFORE
    ${PROJECT_BINARY_DIR}/host/online_compile/include
)

message(STATUS "Hip compiler flags: ${HIP_COMPILER_FLAGS}")

## HIP_COMPILER_FLAGS will be used for on-line compiling of the HIP kernels
set(HIP_COMPILER_FLAGS "${HIP_COMPILER_FLAGS} ${HIP_ONLINE_COMPILER_FLAGS}")
add_definitions("-DHIP_COMPILER_FLAGS=${HIP_COMPILER_FLAGS}")

file(GLOB_RECURSE COMPOSABLE_KERNEL_INCLUDE_1 "${PROJECT_SOURCE_DIR}/composable_kernel/include/*/*.hpp")
file(GLOB COMPOSABLE_KERNEL_INCLUDE_2 "${PROJECT_SOURCE_DIR}/external/rocm/include/bfloat16_dev.hpp")
set(MCONV_KERNEL_INCLUDES
    ${COMPOSABLE_KERNEL_INCLUDE_1}
    ${COMPOSABLE_KERNEL_INCLUDE_2}
   )

file(GLOB_RECURSE MCONV_KERNELS "${PROJECT_SOURCE_DIR}/composable_kernel/src/kernel_wrapper/*.cpp")

add_kernels(${CMAKE_CURRENT_SOURCE_DIR} "${MCONV_KERNELS}")
add_kernel_includes(${CMAKE_CURRENT_SOURCE_DIR} "${MCONV_KERNEL_INCLUDES}")

set(ONLINE_COMPILATION_SOURCE
     ${PROJECT_BINARY_DIR}/kernel.cpp
     ${PROJECT_BINARY_DIR}/kernel_includes.cpp
)

include_directories(BEFORE
    ${PROJECT_BINARY_DIR}/host/online_compile/include
    include
)

set(OLC_HIP_UTILITY_CPPS
    hip_utility/logger.cpp
    hip_utility/tmp_dir.cpp
    hip_utility/md5.cpp  
    hip_utility/exec_utils.cpp
    hip_utility/target_properties.cpp  
    hip_utility/handlehip.cpp
    hip_utility/kernel_build_params.cpp  
    hip_utility/hip_build_utils.cpp  
    hip_utility/hipoc_program.cpp  
    hip_utility/hipoc_kernel.cpp  
    hip_utility/kernel_cache.cpp  
    hip_utility/binary_cache.cpp
   )

list(APPEND OLC_SOURCES ${OLC_HIP_UTILITY_CPPS} ${OLC_HIP_UTILITY_HEADERS})

## addkernels provide the tool to create inlined kernels in one header
add_subdirectory(addkernels)

function(inline_kernels_src KERNELS KERNEL_INCLUDES)
    set(KERNEL_SRC_HPP_FILENAME batch_all.cpp.hpp)
    set(KERNEL_SRC_HPP_PATH ${PROJECT_BINARY_DIR}/inlined_kernels/${KERNEL_SRC_HPP_FILENAME})
    set(KERNEL_SRC_CPP_PATH ${PROJECT_BINARY_DIR}/inlined_kernels/batch_all.cpp)

    add_custom_command(
        OUTPUT ${KERNEL_SRC_HPP_PATH}
        WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}
        DEPENDS addkernels ${KERNELS} ${KERNEL_INCLUDES}
        COMMAND $<TARGET_FILE:addkernels> -target ${KERNEL_SRC_HPP_PATH} -extern -source ${KERNELS}
	COMMENT "Inlining All kernels"
    )
    configure_file(kernels_batch.cpp.in ${KERNEL_SRC_CPP_PATH})
    list(APPEND OLC_SOURCES ${KERNEL_SRC_CPP_PATH} ${KERNEL_SRC_HPP_PATH})

    set(OLC_SOURCES ${OLC_SOURCES} PARENT_SCOPE)
endfunction()

inline_kernels_src("${MCONV_KERNELS}" "${MCONV_KERNEL_INCLUDES}")

list(APPEND ONLINE_COMPILATION_SOURCE ${OLC_SOURCES} ${PROJECT_BINARY_DIR}/olc_kernel_includes.h)

add_custom_command(
    OUTPUT ${PROJECT_BINARY_DIR}/olc_kernel_includes.h
    WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}
    DEPENDS addkernels ${MCONV_KERNEL_INCLUDES}
    COMMAND $<TARGET_FILE:addkernels> -no-recurse -guard GUARD_OLC_KERNEL_INCLUDES_HPP_ -target ${PROJECT_BINARY_DIR}/olc_kernel_includes.h -source ${MCONV_KERNEL_INCLUDES}
    COMMENT "Inlining HIP kernel includes"
  )

## the library target
add_library(online_compile SHARED ${ONLINE_COMPILATION_SOURCE}) 

target_include_directories(online_compile PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/online_compile/include/)
target_include_directories(online_compile PRIVATE ${PROJECT_BINARY_DIR})
target_include_directories(online_compile PRIVATE ${PROJECT_SOURCE_DIR}/external/half/include/)

target_link_libraries(online_compile PRIVATE hip::device)
target_link_libraries(online_compile INTERFACE hip::host)
target_link_libraries(online_compile PRIVATE Boost::filesystem)

target_compile_features(online_compile PUBLIC)
set_target_properties(online_compile PROPERTIES POSITION_INDEPENDENT_CODE ON)

install(TARGETS online_compile LIBRARY DESTINATION lib)