vulkan : cmake integration

.devops/nix/package.nix

@@ -17,6 +17,7 @@
   rocmPackages,
   vulkan-headers,
   vulkan-loader,
+  shaderc,
   clblast,
   useBlas ? builtins.all (x: !x) [
     useCuda
@@ -132,6 +133,8 @@ let
   vulkanBuildInputs = [
     vulkan-headers
     vulkan-loader
+    shaderc
+    python3
   ];
 in
 

.github/workflows/build.yml

@@ -355,8 +355,10 @@ jobs:
       - name: Dependencies
         id: depends
         run: |
-          sudo apt-get update
-          sudo apt-get install build-essential libvulkan-dev
+          wget -qO - https://packages.lunarg.com/lunarg-signing-key-pub.asc | sudo apt-key add -
+          sudo wget -qO /etc/apt/sources.list.d/lunarg-vulkan-jammy.list https://packages.lunarg.com/vulkan/lunarg-vulkan-jammy.list
+          sudo apt-get update -y
+          sudo apt-get install -y build-essential vulkan-sdk
 
       - name: Build
         id: cmake_build

.gitignore

@@ -59,6 +59,8 @@ llama-batched-swift
 /rpc-server
 out/
 tmp/
+ggml-vulkan-shaders.hpp
+ggml-vulkan-shaders.cpp
 
 # CI
 

CMakeLists.txt

@@ -123,7 +123,16 @@ set(LLAMA_INCLUDE_INSTALL_DIR ${CMAKE_INSTALL_INCLUDEDIR} CACHE PATH "Location o
 set(LLAMA_LIB_INSTALL_DIR     ${CMAKE_INSTALL_LIBDIR}     CACHE PATH "Location of library files")
 set(LLAMA_BIN_INSTALL_DIR     ${CMAKE_INSTALL_BINDIR}     CACHE PATH "Location of binary  files")
 
-get_directory_property(LLAMA_TRANSIENT_DEFINES COMPILE_DEFINITIONS)
+
+# At the moment some compile definitions are placed within the ggml/src
+# directory but not exported on the `ggml` target. This could be improved by
+# determining _precisely_ which defines are necessary for the llama-config
+# package.
+#
+get_directory_property(GGML_DIR_DEFINES DIRECTORY ggml/src COMPILE_DEFINITIONS)
+get_target_property(GGML_TARGET_DEFINES ggml COMPILE_DEFINITIONS)
+set(GGML_TRANSIENT_DEFINES ${GGML_TARGET_DEFINES} ${GGML_DIR_DEFINES})
+get_target_property(GGML_LINK_LIBRARIES ggml LINK_LIBRARIES)
 
 set_target_properties(llama PROPERTIES PUBLIC_HEADER ${CMAKE_CURRENT_SOURCE_DIR}/include/llama.h)
 install(TARGETS llama LIBRARY PUBLIC_HEADER)

Makefile

@@ -688,8 +688,8 @@ endif # GGML_CUDA
 
 ifdef GGML_VULKAN
 	MK_CPPFLAGS += -DGGML_USE_VULKAN
-	MK_LDFLAGS  += -lvulkan
-	OBJ_GGML    += ggml/src/ggml-vulkan.o
+	MK_LDFLAGS  += $(shell pkg-config --libs vulkan)
+	OBJ_GGML    += ggml/src/ggml-vulkan.o ggml/src/ggml-vulkan-shaders.o
 
 ifdef GGML_VULKAN_CHECK_RESULTS
 	MK_CPPFLAGS  += -DGGML_VULKAN_CHECK_RESULTS
@@ -711,10 +711,26 @@ ifdef GGML_VULKAN_RUN_TESTS
 	MK_CPPFLAGS  += -DGGML_VULKAN_RUN_TESTS
 endif
 
-ggml/src/ggml-vulkan.o: \
-	ggml/src/ggml-vulkan.cpp \
-	ggml/include/ggml-vulkan.h
+PYTHON_CMD = python
+GLSLC_CMD  = glslc
+_llama_vk_genshaders_cmd = $(PYTHON_CMD) ggml/ggml_vk_generate_shaders.py
+_llama_vk_header = ggml/src/ggml-vulkan-shaders.hpp
+_llama_vk_source = ggml/src/ggml-vulkan-shaders.cpp
+_llama_vk_input_dir = ggml/src/vulkan-shaders
+_llama_vk_shader_deps = $(echo $(_llama_vk_input_dir)/*.comp)
+
+ggml/src/ggml-vulkan.o: ggml/src/ggml-vulkan.cpp ggml/include/ggml-vulkan.h $(_llama_vk_header) $(_llama_vk_source)
 	$(CXX) $(CXXFLAGS) -c $< -o $@
+
+$(_llama_vk_header): $(_llama_vk_source)
+
+$(_llama_vk_source): $(_llama_vk_shader_deps)
+	$(_llama_vk_genshaders_cmd) \
+		--glslc      $(GLSLC_CMD) \
+		--input-dir  $(_llama_vk_input_dir) \
+		--target-hpp $(_llama_vk_header) \
+		--target-cpp $(_llama_vk_source)
+
 endif # GGML_VULKAN
 
 ifdef GGML_HIPBLAS
@@ -1084,6 +1100,7 @@ clean:
 	rm -vrf ggml/src/ggml-cuda/template-instances/*.o
 	rm -rvf $(BUILD_TARGETS)
 	rm -rvf $(TEST_TARGETS)
+	rm -f ggml/src/ggml-vulkan-shaders.hpp ggml/src/ggml-vulkan-shaders.cpp
 	find examples pocs -type f -name "*.o" -delete
 
 #

README.md

@@ -579,6 +579,21 @@ Building the program with BLAS support may lead to some performance improvements
   | GGML_CUDA_KQUANTS_ITER | 1 or 2                 | 2       | Number of values processed per iteration and per HIP thread for Q2_K and Q6_K quantization formats. Setting this value to 1 can improve performance for slow GPUs.                                                                             |
 
 - #### Vulkan
+  ##### Windows (MSYS2)
+  Install [MSYS2](https://www.msys2.org/) and then run the following commands in a UCRT terminal to install dependencies.
+    ```sh
+    pacman -S git \
+        mingw-w64-ucrt-x86_64-gcc \
+        mingw-w64-ucrt-x86_64-cmake \
+        mingw-w64-ucrt-x86_64-vulkan-devel \
+        mingw-w64-ucrt-x86_64-shaderc \
+        mingw-w64-ucrt-x86_64-python3
+    ```
+  Switch into `llama.cpp` directory and build using CMake.
+  ```sh
+  cmake -B build -DLLAMA_VULKAN=ON
+  cmake --build build --config Release
+  ```
 
   **With docker**:
 

cmake/llama-config.cmake.in

@@ -8,6 +8,13 @@ set(GGML_CUDA       @GGML_CUDA@)
 set(GGML_METAL      @GGML_METAL@)
 set(GGML_HIPBLAS    @GGML_HIPBLAS@)
 set(GGML_ACCELERATE @GGML_ACCELERATE@)
+set(GGML_VULKAN @GGML_VULKAN@)
+set(GGML_VULKAN_CHECK_RESULTS @GGML_VULKAN_CHECK_RESULTS@)
+set(GGML_VULKAN_DEBUG @GGML_VULKAN_DEBUG@)
+set(GGML_VULKAN_MEMORY_DEBUG @GGML_VULKAN_MEMORY_DEBUG@)
+set(GGML_VULKAN_VALIDATE @GGML_VULKAN_VALIDATE@)
+set(GGML_SYCL @GGML_SYCL@)
+set(GGML_OPENMP @GGML_OPENMP@)
 
 @PACKAGE_INIT@
 
@@ -37,18 +44,36 @@ if (GGML_METAL)
     find_library(METALKIT_FRAMEWORK MetalKit REQUIRED)
 endif()
 
+if (GGML_VULKAN)
+    find_package(Vulkan REQUIRED)
+endif()
+
 if (GGML_HIPBLAS)
     find_package(hip REQUIRED)
     find_package(hipblas REQUIRED)
     find_package(rocblas REQUIRED)
 endif()
 
+if (GGML_SYCL)
+    find_package(IntelSYCL REQUIRED)
+    find_package(MKL REQUIRED)
+endif()
+
+if (GGML_OPENMP)
+    find_package(OpenMP REQUIRED)
+endif()
+
+
+find_library(ggml_LIBRARY ggml
+    REQUIRED
+    HINTS ${LLAMA_LIB_DIR})
+
 find_library(llama_LIBRARY llama
     REQUIRED
     HINTS ${LLAMA_LIB_DIR})
 
-set(_llama_link_deps "Threads::Threads" "@LLAMA_EXTRA_LIBS@")
-set(_llama_transient_defines "@LLAMA_TRANSIENT_DEFINES@")
+set(_llama_link_deps "${ggml_LIBRARY}" "@GGML_LINK_LIBRARIES@")
+set(_llama_transient_defines "@GGML_TRANSIENT_DEFINES@")
 
 add_library(llama UNKNOWN IMPORTED)
 

ggml/ggml_vk_generate_shaders.py

@@ -5,6 +5,7 @@
 import asyncio
 import os
 from tempfile import gettempdir
+import sys
 
 logger = logging.getLogger("ggml-vk-generate-shaders")
 
@@ -27,9 +28,6 @@
 
 ASYNCIO_CONCURRENCY = 64
 
-input_dir = "vulkan-shaders"
-output_dir = gettempdir()
-
 lock = asyncio.Lock()
 shader_fnames = []
 
@@ -184,30 +182,37 @@ async def withSemaphore(sem, task):
     sem = asyncio.Semaphore(ASYNCIO_CONCURRENCY)
     await asyncio.gather(*(withSemaphore(sem, task) for task in tasks))
 
-    with open("ggml-vulkan-shaders.hpp", "w") as f:
-        f.write("#include <cstdint>\n\n")
+    with open(target_hpp, "w") as hdr, open(target_cpp, "w") as src:
+        hdr.write("#include <cstdint>\n\n")
+        src.write(f"#include \"{os.path.basename(target_hpp)}\"\n\n")
         for name, path in sorted(shader_fnames):
-
             with open(path, "rb") as spv:
                 counter = 0
                 newline_counter = 0
-                f.write(f"unsigned char {name}_data[] = {{\n")
+                data = ""
                 for val in spv.read():
-                    f.write(f"0x{val:02x},")
+                    data += f"0x{val:02x},"
                     newline_counter += 1
                     counter += 1
                     if newline_counter >= 12:
                         newline_counter = 0
-                        f.write("\n")
-            f.write("\n};\n")
-            f.write(f"const uint64_t {name}_len = {counter};\n\n")
-            os.remove(path)
+                        data += "\n"
+            hdr.write(f"extern unsigned char {name}_data[{counter}];\n")
+            hdr.write(f"const uint64_t {name}_len = {counter};\n\n")
+            src.write(f"unsigned char {name}_data[{counter}] = {{\n{data}\n}};\n\n")
+            if not no_clean:
+                os.remove(path)
 
 
 if __name__ == "__main__":
     parser = argparse.ArgumentParser(description="GGML Vulkan Shader Generator")
 
     parser.add_argument("--glslc", help="Path to glslc")
+    parser.add_argument("--input-dir", default="vulkan-shaders", help="Directory containing shader sources")
+    parser.add_argument("--output-dir", default=gettempdir(), help="Directory for containing SPIR-V output")
+    parser.add_argument("--target-hpp", default="ggml-vulkan-shaders.hpp", help="Path to generated header file")
+    parser.add_argument("--target-cpp", default="ggml-vulkan-shaders.cpp", help="Path to generated cpp file")
+    parser.add_argument("--no-clean", action="store_true", help="Keep temporary SPIR-V files in output-dir after build")
     parser.add_argument("--verbose", action="store_true", help="increase output verbosity")
 
     args = parser.parse_args()
@@ -217,4 +222,16 @@ async def withSemaphore(sem, task):
     if args.glslc:
         GLSLC = args.glslc
 
+    input_dir = args.input_dir
+    if not os.path.isdir(input_dir):
+        sys.exit(f"\"{input_dir}\" must be a valid directory containing shader sources")
+
+    output_dir = args.output_dir
+    if not os.path.isdir(output_dir):
+        os.makedirs(output_dir)
+
+    target_hpp = args.target_hpp
+    target_cpp = args.target_cpp
+    no_clean = args.no_clean
+
     asyncio.run(main())

ggml/src/CMakeLists.txt

@@ -525,14 +525,16 @@ if (GGML_RPC)
 endif()
 
 if (GGML_VULKAN)
+    find_package(Python COMPONENTS Interpreter)
+    if (NOT Python_FOUND)
+        message(FATAL_ERROR "python is required to compile Vulkan shaders")
+    endif()
+
     find_package(Vulkan)
 
     if (Vulkan_FOUND)
         message(STATUS "Vulkan found")
 
-        set(GGML_HEADERS_VULKAN ../include/ggml-vulkan.h)
-        set(GGML_SOURCES_VULKAN ggml-vulkan.cpp)
-
         list(APPEND GGML_CDEF_PUBLIC GGML_USE_VULKAN)
 
         # Workaround to the "can't dereference invalidated vector iterator" bug in clang-cl debug build
@@ -561,7 +563,35 @@ if (GGML_VULKAN)
             add_compile_definitions(GGML_VULKAN_RUN_TESTS)
         endif()
 
-        set(GGML_EXTRA_LIBS ${GGML_EXTRA_LIBS} Vulkan::Vulkan)
+        set (_llama_vk_genshaders_cmd ${Python_EXECUTABLE} ${CMAKE_CURRENT_SOURCE_DIR}/../ggml_vk_generate_shaders.py)
+        set (_llama_vk_header     ${CMAKE_CURRENT_BINARY_DIR}/ggml-vulkan-shaders.hpp)
+        set (_llama_vk_source     ${CMAKE_CURRENT_BINARY_DIR}/ggml-vulkan-shaders.cpp)
+        set (_llama_vk_input_dir  ${CMAKE_CURRENT_SOURCE_DIR}/vulkan-shaders)
+        set (_llama_vk_output_dir ${CMAKE_CURRENT_BINARY_DIR}/vulkan-shaders.spv)
+
+        file(GLOB _llama_vk_shader_deps "${_llama_vk_input_dir}/*.comp")
+
+        add_custom_command(
+            OUTPUT ${_llama_vk_header}
+                   ${_llama_vk_source}
+
+            COMMAND ${_llama_vk_genshaders_cmd}
+                --glslc      ${Vulkan_GLSLC_EXECUTABLE}
+                --input-dir  ${_llama_vk_input_dir}
+                --output-dir ${_llama_vk_output_dir}
+                --target-hpp ${_llama_vk_header}
+                --target-cpp ${_llama_vk_source}
+                --no-clean
+
+            DEPENDS ${_llama_vk_shader_deps}
+            COMMENT "Generate vulkan shaders"
+        )
+
+        set(GGML_HEADERS_VULKAN ${CMAKE_CURRENT_SOURCE_DIR}/../include/ggml-vulkan.h ${_llama_vk_header})
+        set(GGML_SOURCES_VULKAN ggml-vulkan.cpp ${_llama_vk_source})
+
+        set(GGML_EXTRA_LIBS     ${GGML_EXTRA_LIBS} Vulkan::Vulkan)
+        set(GGML_EXTRA_INCLUDES ${GGML_EXTRA_INCLUDES} ${CMAKE_CURRENT_BINARY_DIR})
     else()
         message(WARNING "Vulkan not found")
     endif()