Support Moorethreads MUSA

CMakeLists.txt

@@ -32,6 +32,7 @@ option(GGML_NO_ACCELERATE           "ggml: disable Accelerate framework" OFF)
 option(GGML_OPENBLAS                "ggml: use OpenBLAS"                 OFF)
 option(GGML_CLBLAST                 "ggml: use clBLAST"                  OFF)
 option(GGML_CUBLAS                  "ggml: use cuBLAS"                   OFF)
+option(GGML_MUSA                    "ggml: use MUSA"                     OFF)
 option(GGML_METAL                   "ggml: use Metal"                    OFF)
 
 # sanitizers

cmake/CMakeDetermineMUSACompiler.cmake

@@ -0,0 +1,11 @@
+
+set(CMAKE_MUSA_ARCHITECTURES "mp_${MUSA_ARCH}")
+set(CMAKE_MUSA_COMPILER "${MUSA_MCC}")
+set(CMAKE_MUSA_COMPILER_ID "Clang")
+set(CMAKE_MUSA_COMPILER_ARG1 "")
+set(CMAKE_MUSA_COMPILER_ENV_VAR "MCC")
+
+configure_file(
+    ${CMAKE_CURRENT_LIST_DIR}/CMakeMUSACompiler.cmake.in
+    ${CMAKE_PLATFORM_INFO_DIR}/CMakeMUSACompiler.cmake
+)

cmake/CMakeMUSACompiler.cmake.in

@@ -0,0 +1,6 @@
+set(CMAKE_MUSA_COMPILER "@CMAKE_MUSA_COMPILER@")
+set(CMAKE_MUSA_COMPILER_ARG1 "@CMAKE_MUSA_COMPILER_ARG1@")
+set(CMAKE_MUSA_COMPILER_LOADED 1)
+set(CMAKE_MUSA_SOURCE_FILE_EXTENSIONS mu;cu)
+set(CMAKE_MUSA_OUTPUT_EXTENSION .o)
+set(CMAKE_MUSA_COMPILER_ENV_VAR "MUSA")

cmake/CMakeMUSAInformation.cmake

@@ -0,0 +1,26 @@
+
+# reuse cxx things
+
+include(CMakeLanguageInformation)
+include(CMakeCommonLanguageInclude)
+
+include(Compiler/Clang)
+
+__compiler_clang(MUSA)
+__compiler_clang_cxx_standards(MUSA)
+
+set(CMAKE_INCLUDE_FLAG_MUSA "-I")
+
+set(CMAKE_MUSA_RUNTIME_LIBRARY_DEFAULT "SHARED")
+set(CMAKE_MUSA_RUNTIME_LIBRARY_LINK_OPTIONS_STATIC  "")
+set(CMAKE_MUSA_RUNTIME_LIBRARY_LINK_OPTIONS_SHARED  "")
+
+# Populated by CMakeHIPInformation.cmake
+set(CMAKE_MUSA_RUNTIME_LIBRARIES_STATIC "")
+set(CMAKE_MUSA_RUNTIME_LIBRARIES_SHARED "")
+
+# compile a C++ file into an object file
+if(NOT CMAKE_MUSA_COMPILE_OBJECT)
+  set(CMAKE_MUSA_COMPILE_OBJECT
+    "<CMAKE_MUSA_COMPILER> -x musa --cuda-gpu-arch=${CMAKE_MUSA_ARCHITECTURES} -fPIC <DEFINES> <INCLUDES> <FLAGS> -o <OBJECT> -c <SOURCE>")
+endif()

cmake/CMakeTestMUSACompiler.cmake

@@ -0,0 +1 @@
+# do nothing, make cmake happy

cmake/FindMUSA.cmake

@@ -0,0 +1,101 @@
+# find MUSA things
+
+include(${CMAKE_ROOT}/Modules/FindPackageHandleStandardArgs.cmake)
+include(${CMAKE_ROOT}/Modules/SelectLibraryConfigurations.cmake)
+include(${CMAKE_ROOT}/Modules/CMakeFindDependencyMacro.cmake)
+
+if(DEFINED ENV{MUSA_HOME})
+    set(MUSA_HOME $ENV{MUSA_HOME})
+else()
+    set(MUSA_HOME /usr/local/musa)
+endif()
+
+set(MUSA_MCC ${MUSA_HOME}/bin/mcc)
+
+if (DEFINED ENV{MUSA_ARCH})
+    set(MUSA_ARCH $ENV{MUSA_ARCH})
+elseif(NOT MUSA_ARCH)
+    set(MUSA_ARCH "21")
+endif()
+
+if(NOT MUSA_INCLUDE_DIR)
+    set(MUSA_INCLUDE_DIR ${MUSA_HOME}/include)
+endif()
+
+if(NOT MUSA_LIBRARY_DIR)
+    set(MUSA_LIBRARY_DIR ${MUSA_HOME}/lib)
+endif()
+
+if(NOT MUSA_LIBRARIES)
+    find_library(
+        MUSA_MUSA_LIBRARY
+        NAMES libmusa.so
+        PATHS ${MUSA_LIBRARY_DIR}
+    )
+
+    find_library(
+        MUSA_MUBLAS_LIBRARY
+        NAMES libmublas.so
+        PATHS ${MUSA_LIBRARY_DIR}
+    )
+
+    find_library(
+        MUSA_MUSART_LIBRARY
+        NAMES libmusart.so
+        PATHS ${MUSA_LIBRARY_DIR}
+    )
+
+    if(MUSA_MUSA_LIBRARY AND MUSA_MUBLAS_LIBRARY AND MUSA_MUSART_LIBRARY)
+        set(MUSA_LIBRARIES "${MUSA_MUSA_LIBRARY};${MUSA_MUBLAS_LIBRARY};${MUSA_MUSART_LIBRARY}")
+        set(MUSA_MUSA_LIBRARY CACHE STRING "${MUSA_MUSA_LIBRARY}")
+        set(MUSA_MUBLAS_LIBRARY CACHE STRING "${MUSA_MUBLAS_LIBRARY}")
+        set(MUSA_MUSART_LIBRARY CACHE STRING "${MUSA_MUSART_LIBRARY}")
+    endif()
+endif()
+
+if(MUSA_LIBRARIES)
+    if(NOT TARGET MUSA::musa)
+        add_library(MUSA::musa SHARED IMPORTED)
+        set_target_properties(MUSA::musa PROPERTIES
+            IMPORTED_LOCATION ${MUSA_MUSA_LIBRARY}
+            INTERFACE_INCLUDE_DIRECTORIES ${MUSA_INCLUDE_DIR}
+        )
+    endif()
+
+    if(NOT TARGET MUSA::mublas)
+        add_library(MUSA::mublas SHARED IMPORTED)
+        set_target_properties(MUSA::mublas PROPERTIES
+            IMPORTED_LOCATION ${MUSA_MUBLAS_LIBRARY}
+            INTERFACE_INCLUDE_DIRECTORIES ${MUSA_INCLUDE_DIR}
+        )
+    endif()
+
+    if(NOT TARGET MUSA::musart)
+        add_library(MUSA::musart SHARED IMPORTED)
+        set_target_properties(MUSA::musart PROPERTIES
+            IMPORTED_LOCATION ${MUSA_MUSART_LIBRARY}
+            INTERFACE_INCLUDE_DIRECTORIES ${MUSA_INCLUDE_DIR}
+        )
+    endif()
+
+    set(MUSA_INCLUDE_DIR ${MUSA_INCLUDE_DIR} CACHE STRING "")
+    set(MUSA_LIBRARY_DIR ${MUSA_LIBRARY_DIR} CACHE STRING "")
+    set(MUSA_LIBRARIES ${MUSA_LIBRARIES} CACHE STRING "")
+endif()
+
+find_package_handle_standard_args(
+    MUSA
+    REQUIRED_VARS
+    MUSA_ARCH
+    MUSA_MCC
+    MUSA_INCLUDE_DIR
+    MUSA_LIBRARIES
+    MUSA_LIBRARY_DIR
+)
+mark_as_advanced(
+    MUSA_INCLUDE_DIR
+    MUSA_LIBRARIES
+    MUSA_LIBRARY_DIR
+    CMAKE_MUSA_ARCHITECTURES
+    CMAKE_MUSA_COMPILER
+)

src/CMakeLists.txt

@@ -217,6 +217,25 @@ if (GGML_CUBLAS)
     endif()
 endif()
 
+if (GGML_MUSA)
+    option(MUSA_ARCH "MUSA architecture" "21")
+
+    list(APPEND CMAKE_MODULE_PATH "${CMAKE_CURRENT_LIST_DIR}/../cmake")
+
+    find_package(MUSA     REQUIRED)
+
+    message(STATUS "MUSA found")
+
+    enable_language(MUSA)
+
+    add_compile_definitions(GGML_USE_MUSA GGML_USE_CUDA)
+
+    set(GGML_CUDA_SOURCES ggml-cuda.cu ggml-cuda.h)
+    set_source_files_properties(ggml-cuda.cu PROPERTIES LANGUAGE MUSA)
+
+    set(GGML_EXTRA_LIBS ${GGML_EXTRA_LIBS} PUBLIC MUSA::musa MUSA::mublas MUSA::musart)
+endif()
+
 if (GGML_METAL)
     find_library(FOUNDATION_LIBRARY         Foundation              REQUIRED)
     find_library(METAL_FRAMEWORK            Metal                   REQUIRED)

src/ggml-cuda.cu

@@ -64,6 +64,10 @@
 #define cudaStreamWaitEvent(stream, event) hipStreamWaitEvent(stream, event, 0)
 #define cudaStream_t hipStream_t
 #define cudaSuccess hipSuccess
+#elif defined(GGML_USE_MUSA)
+#include <musa.h>
+#include <mublas.h>
+#include "musa_compatible.cuh"
 #else
 #include <cuda_runtime.h>
 #include <cublas_v2.h>