fix compilation after merging the latest develop

cmake/Embed.cmake

@@ -21,10 +21,13 @@
 # OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
 # THE SOFTWARE.
 #####################################################################################
-find_program(EMBED_LD ld)
-find_program(EMBED_OBJCOPY objcopy)
 
-option(EMBED_USE_LD "Use ld to embed data files" OFF)
+option(EMBED_USE_BINARY "Use data file embedding to binary" ON)
+
+if(EMBED_USE_BINARY AND NOT WIN32)
+    find_program(EMBED_LD ld REQUIRED)
+    find_program(EMBED_OBJCOPY objcopy REQUIRED)
+endif()
 
 function(wrap_string)
     set(options)
@@ -53,40 +56,76 @@ function(wrap_string)
     set(${PARSE_VARIABLE} "${lines}" PARENT_SCOPE)
 endfunction()
 
-function(generate_embed_source EMBED_NAME SRC_FILE HEADER_FILE BASE_DIRECTORY)
+function(generate_embed_source EMBED_NAME EMBED_DIR BASE_DIRECTORY)
     set(options)
-    set(oneValueArgs "")
+    set(oneValueArgs)
     set(multiValueArgs SYMBOLS FILES)
     cmake_parse_arguments(PARSE "${options}" "${oneValueArgs}" "${multiValueArgs}" ${ARGN})
 
+    set(RESOURCE_ID 100)
     foreach(SYMBOL FILE IN ZIP_LISTS PARSE_SYMBOLS PARSE_FILES)
+        cmake_path(RELATIVE_PATH FILE BASE_DIRECTORY ${BASE_DIRECTORY} OUTPUT_VARIABLE BASE_NAME)
+        if(EMBED_USE_BINARY AND WIN32)
+            string(TOUPPER "${SYMBOL}" SYMBOL)
+            string(APPEND FILE_IDS "#define IDR_${SYMBOL} ${RESOURCE_ID}\n")
+            string(APPEND RC_MAPPING "IDR_${SYMBOL} TEXTFILE \"${BASE_NAME}\"\n")
+            string(APPEND INIT_KERNELS "        {\"${BASE_NAME}\", resource::read(IDR_${SYMBOL})},\n")
+            math(EXPR RESOURCE_ID "${RESOURCE_ID} + 1" OUTPUT_FORMAT DECIMAL)
+        else()
             set(START_SYMBOL "_binary_${SYMBOL}_start")
             set(LENGTH_SYMBOL "_binary_${SYMBOL}_length")
-        if(EMBED_USE_LD)
+            if(EMBED_USE_BINARY)
                 string(APPEND EXTERNS "
 extern const char ${START_SYMBOL}[];
 extern const size_t _binary_${SYMBOL}_size;
 const auto ${LENGTH_SYMBOL} = reinterpret_cast<size_t>(&_binary_${SYMBOL}_size);
-            ")
+")
             else()
                 string(APPEND EXTERNS "
 extern const char ${START_SYMBOL}[];
 extern const size_t ${LENGTH_SYMBOL};
-            ")
+")
             endif()
-        cmake_path(RELATIVE_PATH FILE BASE_DIRECTORY ${BASE_DIRECTORY} OUTPUT_VARIABLE BASE_NAME)
             string(APPEND INIT_KERNELS "
         { \"${BASE_NAME}\", { ${START_SYMBOL}, ${LENGTH_SYMBOL}} },")
+        endif()
     endforeach()
+    if(EMBED_USE_BINARY AND WIN32)
+       file(WRITE "${EMBED_DIR}/include/resource.h" "
+#define TEXTFILE 256
+
+${FILE_IDS}
+")
+        file(WRITE "${EMBED_DIR}/resource.rc" "
+#include \"resource.h\"
+
+${RC_FILE_MAPPING}
+")
+        set(EXTERNS "
+#include <Windows.h>
+#include \"resource.h\"
 
-    file(WRITE "${HEADER_FILE}" "
+namespace resource {
+std::string_view read(int id)
+{
+    HMODULE handle = GetModuleHandle(nullptr);
+    HRSRC rc = FindResource(handle, MAKEINTRESOURCE(id), MAKEINTRESOURCE(TEXTFILE));
+    HGLOBAL data = LoadResource(handle, rc);
+    return {static_cast<const char*>(LockResource(data)), SizeofResource(handle, rc)};
+}
+}
+
+")
+        set(EMBED_FILES ${EMBED_DIR}/include/resource.h ${EMBED_DIR}/resource.rc)
+    endif()
+    file(WRITE "${EMBED_DIR}/include/${EMBED_NAME}.hpp" "
 #include <string_view>
 #include <unordered_map>
 #include <utility>
 std::unordered_map<std::string_view, std::string_view> ${EMBED_NAME}();
 ")
 
-    file(WRITE "${SRC_FILE}" "
+    file(WRITE "${EMBED_DIR}/${EMBED_NAME}.cpp" "
 #include <${EMBED_NAME}.hpp>
 ${EXTERNS}
 std::unordered_map<std::string_view, std::string_view> ${EMBED_NAME}()
@@ -95,23 +134,28 @@ std::unordered_map<std::string_view, std::string_view> ${EMBED_NAME}()
     return result;
 }
 ")
+    list(APPEND EMBED_FILES ${EMBED_DIR}/${EMBED_NAME}.cpp ${EMBED_DIR}/include/${EMBED_NAME}.hpp)
+    set(EMBED_FILES ${EMBED_FILES} PARENT_SCOPE)
 endfunction()
 
 function(embed_file FILE BASE_DIRECTORY)
     message(STATUS "    ${FILE}")
-    cmake_path(RELATIVE_PATH FILE BASE_DIRECTORY ${BASE_DIRECTORY} OUTPUT_VARIABLE REL_FILE)
+    cmake_path(RELATIVE_PATH FILE BASE_DIRECTORY "${BASE_DIRECTORY}" OUTPUT_VARIABLE REL_FILE)
     string(MAKE_C_IDENTIFIER "${REL_FILE}" OUTPUT_SYMBOL)
     get_filename_component(OUTPUT_FILE_DIR "${REL_FILE}" DIRECTORY)
     file(MAKE_DIRECTORY "${CMAKE_CURRENT_BINARY_DIR}/${OUTPUT_FILE_DIR}")
-    if(EMBED_USE_LD)
+    if(EMBED_USE_BINARY)
+        if(NOT WIN32)
             set(OUTPUT_FILE "${CMAKE_CURRENT_BINARY_DIR}/${REL_FILE}.o")
             add_custom_command(
-            OUTPUT ${OUTPUT_FILE}
+                OUTPUT "${OUTPUT_FILE}"
                 COMMAND ${EMBED_LD} -r -o "${OUTPUT_FILE}" -z noexecstack --format=binary "${REL_FILE}"
                 COMMAND ${EMBED_OBJCOPY} --rename-section .data=.rodata,alloc,load,readonly,data,contents "${OUTPUT_FILE}"
-            WORKING_DIRECTORY ${BASE_DIRECTORY}
-            DEPENDS ${FILE}
+                WORKING_DIRECTORY "${BASE_DIRECTORY}"
+                DEPENDS "${FILE}"
                 VERBATIM)
+            set(OUTPUT_FILE ${OUTPUT_FILE} PARENT_SCOPE)
+        endif()
     else()
         set(OUTPUT_FILE "${CMAKE_CURRENT_BINARY_DIR}/${REL_FILE}.cpp")
         # reads source file contents as hex string
@@ -127,40 +171,38 @@ function(embed_file FILE BASE_DIRECTORY)
 extern const char _binary_${OUTPUT_SYMBOL}_start[] = { ${ARRAY_VALUES} };
 extern const size_t _binary_${OUTPUT_SYMBOL}_length = sizeof(_binary_${OUTPUT_SYMBOL}_start);
 ")
-    endif()
         set(OUTPUT_FILE ${OUTPUT_FILE} PARENT_SCOPE)
+    endif()
     set(OUTPUT_SYMBOL ${OUTPUT_SYMBOL} PARENT_SCOPE)
 endfunction()
 
 function(add_embed_library EMBED_NAME)
     set(options)
-    set(oneValueArgs BASE_DIRECTORY)
+    set(oneValueArgs RELATIVE)
     set(multiValueArgs)
     cmake_parse_arguments(PARSE "${options}" "${oneValueArgs}" "${multiValueArgs}" ${ARGN})
+
     set(EMBED_DIR ${CMAKE_CURRENT_BINARY_DIR}/embed/${EMBED_NAME})
     file(MAKE_DIRECTORY ${EMBED_DIR})
-    set(SRC_FILE "${EMBED_DIR}/${EMBED_NAME}.cpp")
-    set(HEADER_FILE "${EMBED_DIR}/include/${EMBED_NAME}.hpp")
     message(STATUS "Embedding kernel files:")
     foreach(FILE ${PARSE_UNPARSED_ARGUMENTS})
-        embed_file(${FILE} ${PARSE_BASE_DIRECTORY})
+        embed_file(${FILE} ${PARSE_RELATIVE})
         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_FILE} ${HEADER_FILE} "${PARSE_BASE_DIRECTORY}" SYMBOLS ${SYMBOLS} FILES ${PARSE_UNPARSED_ARGUMENTS})
-    add_library(embed_lib_${EMBED_NAME} OBJECT ${SRC_FILE} ${HEADER_FILE})
-    if(NOT EMBED_USE_LD)
-        target_sources(embed_lib_${EMBED_NAME} PRIVATE ${OUTPUT_FILES})
+    generate_embed_source(${EMBED_NAME} ${EMBED_DIR} "${PARSE_RELATIVE}" SYMBOLS ${SYMBOLS} FILES ${PARSE_UNPARSED_ARGUMENTS})
+    set(INTERNAL_EMBED_LIB embed_lib_${EMBED_NAME})
+    add_library(${INTERNAL_EMBED_LIB} OBJECT ${EMBED_FILES})
+    if(NOT EMBED_USE_BINARY)
+        target_sources(${INTERNAL_EMBED_LIB} PRIVATE ${OUTPUT_FILES})
     endif()
-    target_include_directories(embed_lib_${EMBED_NAME} PUBLIC ${EMBED_DIR}/include)
-    target_compile_options(embed_lib_${EMBED_NAME} PRIVATE
-            -Wno-reserved-identifier -Wno-extern-initializer -Wno-missing-variable-declarations)
-    set_target_properties(embed_lib_${EMBED_NAME} PROPERTIES POSITION_INDEPENDENT_CODE ON)
-    add_library(${EMBED_NAME} INTERFACE $<TARGET_OBJECTS:embed_lib_${EMBED_NAME}> ${OUTPUT_FILES})
-    target_link_libraries(${EMBED_NAME} INTERFACE $<TARGET_OBJECTS:embed_lib_${EMBED_NAME}>)
-    if(EMBED_USE_LD)
-        target_link_libraries(${EMBED_NAME} INTERFACE ${OUTPUT_FILES})
+    target_include_directories(${INTERNAL_EMBED_LIB} PRIVATE "${EMBED_DIR}/include")
+    target_compile_options(${INTERNAL_EMBED_LIB} PRIVATE -Wno-reserved-identifier -Wno-extern-initializer -Wno-missing-variable-declarations)
+    set_target_properties(${INTERNAL_EMBED_LIB} PROPERTIES POSITION_INDEPENDENT_CODE On)
+    add_library(${EMBED_NAME} INTERFACE $<TARGET_OBJECTS:${INTERNAL_EMBED_LIB}> ${OUTPUT_FILES})
+    if(EMBED_USE_BINARY AND WIN32)
+        target_link_libraries(${EMBED_NAME} INTERFACE $<TARGET_OBJECTS:${INTERNAL_EMBED_LIB}>)
     endif()
-    target_include_directories(${EMBED_NAME} INTERFACE ${EMBED_DIR}/include)
+    target_include_directories(${EMBED_NAME} INTERFACE "${EMBED_DIR}/include")
 endfunction()

extern/CMakeLists.txt

@@ -188,11 +188,11 @@ ExternalProject_Add(
         BUILD_COMMAND ${NMAKE_EXECUTABLE} /f ..\\sqlite3\\Makefile.msc USE_AMALGAMATION=1 NO_TCL=1 TOP=..\\sqlite3 libsqlite3.lib
         INSTALL_COMMAND "")
 
+add_library(SQLite::SQLite3 INTERFACE IMPORTED GLOBAL)
+add_dependencies(SQLite::SQLite3 sqlite3)
+
 ExternalProject_Get_Property(sqlite3 BINARY_DIR)
 
-# For compatibility with PkgConfig on Linux
-add_library(PkgConfig::SQLITE3 INTERFACE IMPORTED GLOBAL)
-add_dependencies(PkgConfig::SQLITE3 sqlite3)
-target_link_directories(PkgConfig::SQLITE3 INTERFACE ${BINARY_DIR})
-target_link_libraries(PkgConfig::SQLITE3 INTERFACE libsqlite3.lib)
-target_include_directories(PkgConfig::SQLITE3 INTERFACE ${BINARY_DIR})
+target_link_directories(SQLite::SQLite3 INTERFACE ${BINARY_DIR})
+target_link_libraries(SQLite::SQLite3 INTERFACE libsqlite3.lib)
+target_include_directories(SQLite::SQLite3 INTERFACE ${BINARY_DIR})

src/targets/gpu/CMakeLists.txt

@@ -54,10 +54,8 @@ if(NOT MIGRAPHX_USE_COMPOSABLEKERNEL)
         ${CMAKE_CURRENT_SOURCE_DIR}/kernels/include/migraphx/kernels/ck.hpp)
 endif()
 
-if(NOT WIN32)
-    include(Embed)
-    add_embed_library(migraphx_kernels ${KERNEL_FILES} BASE_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/kernels/include/)
-endif()
+include(Embed)
+add_embed_library(migraphx_kernels ${KERNEL_FILES} RELATIVE ${CMAKE_CURRENT_SOURCE_DIR}/kernels/include/)
 
 configure_file(device/targets.hpp.in include/migraphx/gpu/device/targets.hpp)
 file(GLOB DEVICE_GPU_SRCS CONFIGURE_DEPENDS ${CMAKE_CURRENT_SOURCE_DIR}/device/*.cpp)
@@ -158,20 +156,6 @@ add_library(migraphx_gpu
 set_target_properties(migraphx_gpu PROPERTIES EXPORT_NAME gpu)
 migraphx_generate_export_header(migraphx_gpu)
 
-if(WIN32)
-    target_include_directories(migraphx_gpu PRIVATE kernels)
-    target_sources(migraphx_gpu
-        PRIVATE
-            kernels/migraphx_kernels.cpp
-            kernels/migraphx_kernels.hpp
-            kernels/resource.h
-            kernels/resource.rc)
-    target_compile_options(migraphx_gpu PUBLIC $<BUILD_INTERFACE:-Wno-ignored-attributes>)
-else()
-    target_link_libraries(migraphx_gpu PRIVATE migraphx_kernels)
-    target_link_options(migraphx_gpu PUBLIC -Wno-option-ignored)
-endif()
-
 function(register_migraphx_gpu_ops PREFIX)
     foreach(OP ${ARGN})
         register_op(migraphx_gpu HEADER migraphx/gpu/${OP}.hpp OPERATORS gpu::${PREFIX}${OP} INCLUDES migraphx/gpu/context.hpp)
@@ -295,7 +279,7 @@ else()
 endif()
 
 target_link_libraries(migraphx_gpu PUBLIC migraphx MIOpen roc::rocblas)
-target_link_libraries(migraphx_gpu PRIVATE migraphx_device)
+target_link_libraries(migraphx_gpu PRIVATE migraphx_device migraphx_kernels)
 if(MIGRAPHX_USE_COMPOSABLEKERNEL)
     target_link_libraries(migraphx_gpu PRIVATE composable_kernel::jit_library)
 endif()

test/CMakeLists.txt

@@ -42,18 +42,8 @@ foreach(TEST ${TESTS})
 endforeach()
 
 if(MIGRAPHX_ENABLE_GPU)
-    set(GPU_TESTS
-        gpu/adjust_allocation.cpp
-        gpu/context_serialize.cpp
-        gpu/hip.cpp
-        gpu/jit.cpp
-        gpu/literal.cpp
-        gpu/manage_host_buffer.cpp
-        gpu/mlir.cpp
-        gpu/pack_args.cpp
-        gpu/pack_int8_args.cpp
-        gpu/quantization.cpp
-        gpu/stream_sync.cpp)
+    # gpu tests
+    file(GLOB GPU_TESTS CONFIGURE_DEPENDS gpu/*.cpp)
 
     foreach(TEST ${GPU_TESTS})
         get_filename_component(BASE_NAME ${TEST} NAME_WE)
@@ -63,22 +53,16 @@ if(MIGRAPHX_ENABLE_GPU)
             COST 10
             RESOURCE_LOCK gpu
         )
-        if(WIN32)
-            target_include_directories(test_gpu_${BASE_NAME} PRIVATE ${CMAKE_SOURCE_DIR}/src/targets/gpu/kernels)
-        else()
-            target_link_libraries(test_gpu_${BASE_NAME} migraphx_kernels)
-        endif()
         if(MIGRAPHX_USE_HIPRTC)
         target_compile_definitions(test_gpu_${BASE_NAME} PUBLIC -DMIGRAPHX_USE_HIPRTC)
         endif()
-        target_link_libraries(test_gpu_${BASE_NAME} migraphx_gpu)
+        target_link_libraries(test_gpu_${BASE_NAME} migraphx_gpu migraphx_kernels)
     endforeach()
 endif()
 
 if(MIGRAPHX_ENABLE_FPGA)
-    set(FPGA_TESTS
-            fpga/get_target_assignments.cpp
-            fpga/test_compile.cpp)
+    # fpga tests
+    file(GLOB FPGA_TESTS CONFIGURE_DEPENDS fpga/*.cpp)
 
     foreach(TEST ${FPGA_TESTS})
         get_filename_component(BASE_NAME ${TEST} NAME_WE)
@@ -94,10 +78,7 @@ endif()
 
 # Onnx test
 set(TEST_ONNX_DIR ${CMAKE_CURRENT_SOURCE_DIR}/onnx)
-set(ONNX_TESTS
-    onnx/onnx_rnn_test.cpp
-    onnx/onnx_test.cpp
-    onnx/verify_onnx.cpp)
+file(GLOB ONNX_TESTS ${TEST_ONNX_DIR}/*.cpp)
 
 foreach(ONNX_TEST ${ONNX_TESTS})
     get_filename_component(BASE_NAME ${ONNX_TEST} NAME_WE)
@@ -164,36 +145,31 @@ int main() {}\n"
 endfunction()
 
 function(test_headers PREFIX)
-    cmake_parse_arguments(PARSE "" "" "HEADERS;DEPENDS" ${ARGN})
-    foreach(HEADER ${PARSE_HEADERS})
+    file(GLOB HEADERS CONFIGURE_DEPENDS ${ARGN})
+    if(NOT MIGRAPHX_USE_COMPOSABLEKERNEL)
+        list(REMOVE_ITEM HEADERS
+              ${CMAKE_SOURCE_DIR}/src/targets/gpu/include/migraphx/gpu/ck.hpp)
+    endif()
+    foreach(HEADER ${HEADERS})
         file(RELATIVE_PATH HEADER_REL ${CMAKE_SOURCE_DIR} ${HEADER})
         string(MAKE_C_IDENTIFIER ${HEADER_REL} TEST_NAME)
         get_filename_component(BASE_NAME ${HEADER} NAME_WE)
         test_header(header_${TEST_NAME} ${PREFIX}/${BASE_NAME}.hpp)
-        target_link_libraries(header_${TEST_NAME} migraphx migraphx_onnx migraphx_tf migraphx_all_targets ${PARSE_DEPENDS})
+        target_link_libraries(header_${TEST_NAME} migraphx migraphx_onnx migraphx_tf migraphx_all_targets)
     endforeach()
 endfunction()
 
-file(GLOB HEADERS CONFIGURE_DEPENDS ${CMAKE_SOURCE_DIR}/src/include/migraphx/*.hpp)
-test_headers(migraphx HEADERS ${HEADERS})
-
-file(GLOB HEADERS CONFIGURE_DEPENDS ${CMAKE_SOURCE_DIR}/src/targets/ref/include/migraphx/ref/*.hpp)
-test_headers(migraphx/ref HEADERS ${HEADERS})
+test_headers(migraphx ${CMAKE_SOURCE_DIR}/src/include/migraphx/*.hpp)
+test_headers(migraphx/ref ${CMAKE_SOURCE_DIR}/src/targets/ref/include/migraphx/ref/*.hpp)
 
 if(MIGRAPHX_ENABLE_GPU)
-    file(GLOB HEADERS CONFIGURE_DEPENDS ${CMAKE_SOURCE_DIR}/src/targets/gpu/include/migraphx/gpu/*.hpp)
-    if(NOT MIGRAPHX_USE_COMPOSABLEKERNEL)
-        list(REMOVE_ITEM HEADERS ${CMAKE_SOURCE_DIR}/src/targets/gpu/include/migraphx/gpu/ck.hpp)
-    endif()
-    test_headers(migraphx/gpu HEADERS ${HEADERS} DEPENDS migraphx_gpu)
+    test_headers(migraphx/gpu HEADERS ${CMAKE_SOURCE_DIR}/src/targets/gpu/include/migraphx/gpu/*.hpp DEPENDS migraphx_gpu)
 endif()
 if(MIGRAPHX_ENABLE_CPU)
-    file(GLOB HEADERS CONFIGURE_DEPENDS ${CMAKE_SOURCE_DIR}/src/targets/gpu/include/migraphx/cpu/*.hpp)
-    test_headers(migraphx/cpu ${HEADERS} DEPENDS migraphx_cpu)
+    test_headers(migraphx/cpu HEADERS ${CMAKE_SOURCE_DIR}/src/targets/cpu/include/migraphx/cpu/*.hpp migraphx_cpu)
 endif()
 if(MIGRAPHX_ENABLE_FPGA)
-    file(GLOB HEADERS CONFIGURE_DEPENDS ${CMAKE_SOURCE_DIR}/src/targets/gpu/include/migraphx/fpga/*.hpp)
-    test_headers(migraphx/fpga HEADERS ${HEADERS} DEPENDS migraphx_fpga)
+    test_headers(migraphx/fpga HEADERS ${CMAKE_SOURCE_DIR}/src/targets/fpga/include/migraphx/fpga/*.hpp migraphx_fpga)
 endif()