mirror of
https://github.com/ROCm/composable_kernel.git
synced 2026-05-14 10:09:41 +00:00
817b2a47c6
* Add online-compiling facility
* Synchronize from fwd-v4r5 and implement host interfaces to call conv-fwd v4r4/v4r5 using on-line compiling method
* Tiny adjustment to time reporting
* Use object assignment to replace explicit bytes copying in the first kernel of v4r4/v4r5
* Use single thread to assign descriptor object to device memory
* Adjust to the workload assignment of the two kernels of v4r4 (experimental)
* Revert "Adjust to the workload assignment of the two kernels of v4r4 (experimental)"
This reverts commit eb38461456bb0c82b6c0d32cdd616e181907e20c.
* Update to make constexpr for generating descriptor types in kernel 2 of dynamic conv-fwd v4r4
* Update to dynamic conv-fwd v4r4 online-compiling
* Update to dynamic conv-fwd v4r5 online-compiling (result not accurate)
* Tiny update to driver/CMakeLists.txt
* clang-format
* Tiny comments change
* Add env OLC_DUMP_SAVE_TMP_DIR to support saving of temperary dir
* Fwd v4r5 olc perf (#39)
* added hip-clang flags that fix perf issue of online compilation
* fix bug for olc fwd-v4r5-nchw
* Move constexpr and type reference statements out of the function body in conv-fwd v4r4/v4r5 kernel wrapper
* Remove printing in hip_build_utils.cpp
* Update to root CMakeLists.txt
* Revert "Move constexpr and type reference statements out of the function body in conv-fwd v4r4/v4r5 kernel wrapper"
This reverts commit 3d2c5d8ecdd8298b72d127110500ed5b38d9835c.
Co-authored-by: Chao Liu <chao.liu2@amd.com>
Co-authored-by: Chao Liu <lc.roy86@gmail.com>
Co-authored-by: root <root@dc-smc-18.amd.com>
[ROCm/composable_kernel commit: 1685048a67]
41 lines
1.9 KiB
CMake
41 lines
1.9 KiB
CMake
|
|
function(add_kernels SRC_DIR KERNEL_FILES)
|
|
set(INIT_KERNELS_LIST)
|
|
set(KERNELS_DECLS)
|
|
foreach(KERNEL_FILE ${KERNEL_FILES})
|
|
if("${CMAKE_VERSION}" VERSION_LESS 3.0)
|
|
configure_file(${KERNEL_FILE} ${KERNEL_FILE}.delete)
|
|
else()
|
|
set_property(DIRECTORY APPEND PROPERTY CMAKE_CONFIGURE_DEPENDS ${KERNEL_FILE})
|
|
endif()
|
|
get_filename_component(BASE_NAME ${KERNEL_FILE} NAME_WE)
|
|
string(TOUPPER "${BASE_NAME}" KEY_NAME)
|
|
string(MAKE_C_IDENTIFIER "${KEY_NAME}" VAR_NAME)
|
|
string(APPEND KERNELS_DECLS "extern const size_t APP_KERNEL_${VAR_NAME}_SIZE;\n")
|
|
string(APPEND KERNELS_DECLS "extern const unsigned char APP_KERNEL_${VAR_NAME}[];\n")
|
|
list(APPEND INIT_KERNELS_LIST " { \"${KEY_NAME}\", std::string(reinterpret_cast<const char*>(APP_KERNEL_${VAR_NAME}), APP_KERNEL_${VAR_NAME}_SIZE) }")
|
|
endforeach()
|
|
string(REPLACE ";" ",\n" INIT_KERNELS "${INIT_KERNELS_LIST}")
|
|
configure_file(${SRC_DIR}/kernel.cpp.in ${PROJECT_BINARY_DIR}/kernel.cpp)
|
|
endfunction()
|
|
|
|
function(add_kernel_includes SRC_DIR KERNEL_FILES)
|
|
set(INIT_KERNELS_LIST)
|
|
foreach(KERNEL_FILE ${KERNEL_FILES})
|
|
if("${CMAKE_VERSION}" VERSION_LESS 3.0)
|
|
configure_file(${KERNEL_FILE} ${KERNEL_FILE}.delete)
|
|
else()
|
|
set_property(DIRECTORY APPEND PROPERTY CMAKE_CONFIGURE_DEPENDS ${KERNEL_FILE})
|
|
endif()
|
|
get_filename_component(BASE_NAME ${KERNEL_FILE} NAME_WE)
|
|
get_filename_component(FILE_NAME ${KERNEL_FILE} NAME)
|
|
string(TOUPPER "${BASE_NAME}" KEY_NAME)
|
|
string(MAKE_C_IDENTIFIER "${KEY_NAME}" VAR_NAME)
|
|
list(APPEND INIT_KERNELS_LIST " { \"${FILE_NAME}\", std::string(reinterpret_cast<const char*>(${VAR_NAME}), ${VAR_NAME}_SIZE) }")
|
|
endforeach()
|
|
string(REPLACE ";" ",\n" INIT_KERNELS "${INIT_KERNELS_LIST}")
|
|
configure_file(${SRC_DIR}/kernel_includes.cpp.in ${PROJECT_BINARY_DIR}/kernel_includes.cpp)
|
|
endfunction()
|
|
|
|
|