rocm: Automatically build with separate compiler

CMakeLists.txt

@@ -372,41 +372,63 @@ if (LLAMA_CLBLAST)
 endif()
 
 if (LLAMA_HIPBLAS)
-    list(APPEND CMAKE_PREFIX_PATH /opt/rocm)
+    if (WIN32)
+        # todo: also allow building with a separate compiler on windows
+        list(APPEND CMAKE_PREFIX_PATH /opt/rocm)
 
-    if (NOT ${CMAKE_C_COMPILER_ID} MATCHES "Clang")
-        message(WARNING "Only LLVM is supported for HIP, hint: CC=/opt/rocm/llvm/bin/clang")
-    endif()
-    if (NOT ${CMAKE_CXX_COMPILER_ID} MATCHES "Clang")
-        message(WARNING "Only LLVM is supported for HIP, hint: CXX=/opt/rocm/llvm/bin/clang++")
-    endif()
-
-    find_package(hip)
-    find_package(hipblas)
-    find_package(rocblas)
-
-    if (${hipblas_FOUND} AND ${hip_FOUND})
-        message(STATUS "HIP and hipBLAS found")
-        add_compile_definitions(GGML_USE_HIPBLAS GGML_USE_CUBLAS)
-        add_library(ggml-rocm OBJECT ggml-cuda.cu ggml-cuda.h)
-        if (BUILD_SHARED_LIBS)
-            set_target_properties(ggml-rocm PROPERTIES POSITION_INDEPENDENT_CODE ON)
+        if (NOT ${CMAKE_C_COMPILER_ID} MATCHES "Clang")
+            message(WARNING "Only LLVM is supported for HIP, hint: CC=/opt/rocm/llvm/bin/clang")
         endif()
-        if (LLAMA_CUDA_FORCE_DMMV)
-            target_compile_definitions(ggml-rocm PRIVATE GGML_CUDA_FORCE_DMMV)
+        if (NOT ${CMAKE_CXX_COMPILER_ID} MATCHES "Clang")
+            message(WARNING "Only LLVM is supported for HIP, hint: CXX=/opt/rocm/llvm/bin/clang++")
         endif()
-        target_compile_definitions(ggml-rocm PRIVATE GGML_CUDA_DMMV_X=${LLAMA_CUDA_DMMV_X})
-        target_compile_definitions(ggml-rocm PRIVATE GGML_CUDA_MMV_Y=${LLAMA_CUDA_MMV_Y})
-        target_compile_definitions(ggml-rocm PRIVATE K_QUANTS_PER_ITERATION=${LLAMA_CUDA_KQUANTS_ITER})
-        set_source_files_properties(ggml-cuda.cu PROPERTIES LANGUAGE CXX)
-        target_link_libraries(ggml-rocm PRIVATE hip::device PUBLIC hip::host roc::rocblas roc::hipblas)
 
+        find_package(hip)
+        find_package(hipblas)
+        find_package(rocblas)
+
+        if (${hipblas_FOUND} AND ${hip_FOUND})
+            message(STATUS "HIP and hipBLAS found")
+            add_compile_definitions(GGML_USE_HIPBLAS GGML_USE_CUBLAS)
+            add_library(ggml-rocm OBJECT ggml-cuda.cu ggml-cuda.h)
+            if (LLAMA_CUDA_FORCE_DMMV)
+                target_compile_definitions(ggml-rocm PRIVATE GGML_CUDA_FORCE_DMMV)
+            endif()
+            target_compile_definitions(ggml-rocm PRIVATE GGML_CUDA_DMMV_X=${LLAMA_CUDA_DMMV_X})
+            target_compile_definitions(ggml-rocm PRIVATE GGML_CUDA_MMV_Y=${LLAMA_CUDA_MMV_Y})
+            target_compile_definitions(ggml-rocm PRIVATE K_QUANTS_PER_ITERATION=${LLAMA_CUDA_KQUANTS_ITER})
+            target_compile_definitions(ggml-rocm PRIVATE CC_TURING=1000000000)
+            set_source_files_properties(ggml-cuda.cu PROPERTIES LANGUAGE CXX)
+            target_link_libraries(ggml-rocm PRIVATE hip::device PUBLIC hip::host roc::rocblas roc::hipblas)
+
+            if (LLAMA_STATIC)
+                message(FATAL_ERROR "Static linking not supported for HIP/ROCm")
+            endif()
+            set(LLAMA_EXTRA_LIBS ${LLAMA_EXTRA_LIBS} ggml-rocm)
+        else()
+            message(WARNING "hipBLAS or HIP not found. Try setting CMAKE_PREFIX_PATH=/opt/rocm")
+        endif()
+    else()
         if (LLAMA_STATIC)
             message(FATAL_ERROR "Static linking not supported for HIP/ROCm")
         endif()
+        include(ExternalProject)
+        if (NOT DEFINED LLAMA_ROCM_PROJECT_CMAKE_ARGS)
+            set(LLAMA_ROCM_PROJECT_CMAKE_ARGS -DCMAKE_CXX_COMPILER=/opt/rocm/llvm/bin/clang++ -DCMAKE_BUILD_TYPE=${CMAKE_BUILD_TYPE})
+        endif()
+        ExternalProject_Add(ggml-rocm-project
+            SOURCE_DIR ../ggml-rocm
+            PREFIX ggml-rocm
+            INSTALL_COMMAND ""
+            CMAKE_ARGS ${LLAMA_ROCM_PROJECT_CMAKE_ARGS}
+            BUILD_BYPRODUCTS ${CMAKE_BINARY_DIR}/bin/libggml-rocm.so
+        )
+        add_compile_definitions(GGML_USE_CUBLAS)
+        add_library(ggml-rocm SHARED IMPORTED)
+        add_dependencies(ggml-rocm ggml-rocm-project)
+        set_target_properties(ggml-rocm PROPERTIES IMPORTED_LOCATION ${CMAKE_BINARY_DIR}/bin/libggml-rocm.so)
         set(LLAMA_EXTRA_LIBS ${LLAMA_EXTRA_LIBS} ggml-rocm)
-    else()
-        message(WARNING "hipBLAS or HIP not found. Try setting CMAKE_PREFIX_PATH=/opt/rocm")
+        set(GGML_SOURCES_EXTRA ${GGML_SOURCES_EXTRA} ggml-cuda.h)
     endif()
 endif()
 
@@ -762,6 +784,14 @@ if (LLAMA_METAL)
             WORLD_READ
         DESTINATION ${CMAKE_INSTALL_BINDIR})
 endif()
+install(
+    FILES ${CMAKE_BINARY_DIR}/../ggml-rocm/src/ggml-rocm-project-build/libggml-rocm.so
+    PERMISSIONS
+        OWNER_READ
+        OWNER_WRITE
+        GROUP_READ
+        WORLD_READ
+    DESTINATION ${CMAKE_INSTALL_BINDIR})
 
 #
 # programs, examples and tests

ggml-rocm/CMakeLists.txt

@@ -0,0 +1,57 @@
+cmake_minimum_required(VERSION 3.12) # Don't bump this version for no reason
+project(ggml-rocm CXX)
+
+set(LLAMA_CUDA_DMMV_X      "32" CACHE STRING "llama: x stride for dmmv CUDA kernels")
+set(LLAMA_CUDA_MMV_Y        "1" CACHE STRING "llama: y block size for mmv CUDA kernels")
+set(LLAMA_CUDA_KQUANTS_ITER "2" CACHE STRING "llama: iters./thread per block for Q2_K/Q6_K")
+
+set(CMAKE_CXX_STANDARD 11)
+set(CMAKE_CXX_STANDARD_REQUIRED true)
+set(THREADS_PREFER_PTHREAD_FLAG ON)
+find_package(Threads REQUIRED)
+
+list(APPEND CMAKE_PREFIX_PATH /opt/rocm)
+
+if (NOT ${CMAKE_CXX_COMPILER_ID} MATCHES "Clang")
+    message(FATAL_ERROR "Only LLVM is supported for HIP, hint: CXX=/opt/rocm/llvm/bin/clang++")
+endif()
+
+find_package(hip)
+find_package(hipblas)
+find_package(rocblas)
+
+if (${hipblas_FOUND} AND ${hip_FOUND})
+    message(STATUS "HIP and hipBLAS found")
+    add_compile_definitions(GGML_USE_HIPBLAS GGML_USE_CUBLAS)
+    add_library(ggml-rocm SHARED ggml-cuda.cu ggml-cuda.h ggml.h)
+    set_target_properties(ggml-rocm PROPERTIES POSITION_INDEPENDENT_CODE ON)
+    if (LLAMA_CUDA_FORCE_DMMV)
+        target_compile_definitions(ggml-rocm PRIVATE GGML_CUDA_FORCE_DMMV)
+    endif()
+    target_compile_definitions(ggml-rocm PRIVATE GGML_CUDA_DMMV_X=${LLAMA_CUDA_DMMV_X})
+    target_compile_definitions(ggml-rocm PRIVATE GGML_CUDA_MMV_Y=${LLAMA_CUDA_MMV_Y})
+    target_compile_definitions(ggml-rocm PRIVATE K_QUANTS_PER_ITERATION=${LLAMA_CUDA_KQUANTS_ITER})
+    target_compile_definitions(ggml-rocm PRIVATE CC_TURING=1000000000)
+    set_source_files_properties(ggml-cuda.cu PROPERTIES LANGUAGE CXX)
+    set_target_properties(ggml-rocm PROPERTIES LINKER_LANGUAGE CXX LIBRARY_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR}/../../../bin)
+    target_link_libraries(ggml-rocm PUBLIC hip::device hip::host roc::rocblas roc::hipblas)
+else()
+    message(FATAL_ERROR "hipBLAS or HIP not found. Try setting CMAKE_PREFIX_PATH=/opt/rocm")
+endif()
+
+if (NOT MSVC)
+    set(cxx_flags
+            -Wall
+            -Wextra
+            -Wpedantic
+            -Wcast-qual
+            -Wno-unused-function
+            -Wno-multichar
+    )
+else()
+    # todo : msvc
+endif()
+
+add_compile_options(
+        "$<$<COMPILE_LANGUAGE:CXX>:${cxx_flags}>"
+)

ggml-rocm/ggml-cuda.cu

@@ -0,0 +1 @@
+../ggml-cuda.cu

ggml-rocm/ggml-cuda.h

@@ -0,0 +1 @@
+../ggml-cuda.h

ggml-rocm/ggml.h

@@ -0,0 +1 @@
+../ggml.h

tests/CMakeLists.txt

@@ -40,4 +40,4 @@ llama_build_and_test_executable(test-grad0.cpp) # SLOW
 # dummy executable - not installed
 get_filename_component(TEST_TARGET test-c.c NAME_WE)
 add_executable(${TEST_TARGET} test-c.c)
-target_link_libraries(${TEST_TARGET} PRIVATE llama)
+target_link_libraries(${TEST_TARGET} PRIVATE llama ggml)