Nouveau Backend (en cours)

- utilise les wrapper c++
- teste reference sur CPU (fp32 non optimisé)
- init de l' iGPU RDNA3
- allocation memoire sur ce backend HIP

- 1ere version du produit matriciel wmma

Author: Djip007

ggml/CMakeLists.txt

@@ -141,6 +141,8 @@ set(GGML_BLAS_VENDOR ${GGML_BLAS_VENDOR_DEFAULT} CACHE STRING
                                             "ggml: BLAS library vendor")
 option(GGML_LLAMAFILE                       "ggml: use LLAMAFILE"                             ${GGML_LLAMAFILE_DEFAULT})
 
+option(GGML_IGPU                            "ggml: use IGPU"                                  OFF)
+
 option(GGML_CUDA                            "ggml: use CUDA"                                  OFF)
 option(GGML_MUSA                            "ggml: use MUSA"                                  OFF)
 option(GGML_CUDA_FORCE_MMQ                  "ggml: use mmq kernels instead of cuBLAS"         OFF)
@@ -244,6 +246,7 @@ set(GGML_PUBLIC_HEADERS
     include/ggml-blas.h
     include/ggml-cann.h
     include/ggml-cuda.h
+    include/ggml-igpu.h
     include/ggml-kompute.h
     include/ggml-opt.h
     include/ggml-metal.h

ggml/cmake/ggml-config.cmake.in

@@ -76,6 +76,11 @@ if (NOT GGML_SHARED_LIB)
         list(APPEND GGML_HIP_INTERFACE_LINK_LIBRARIES hip::host roc::rocblas roc::hipblas)
     endif()
 
+    if (GGML_IGPU)
+        find_package(hip     REQUIRED)
+        list(APPEND GGML_HIP_INTERFACE_LINK_LIBRARIES hip::host)
+    endif()
+
     if (GGML_SYCL)
         find_package(DNNL)
         if (${DNNL_FOUND} AND GGML_SYCL_TARGET STREQUAL "INTEL")

ggml/include/ggml-igpu.h

@@ -0,0 +1,16 @@
+#pragma once
+
+#include "ggml.h"
+#include "ggml-backend.h"
+
+
+#ifdef  __cplusplus
+extern "C" {
+#endif
+
+GGML_BACKEND_API ggml_backend_reg_t ggml_backend_igpu_reg(void);
+
+
+#ifdef  __cplusplus
+}
+#endif

ggml/src/CMakeLists.txt

@@ -324,6 +324,7 @@ ggml_add_backend(BLAS)
 ggml_add_backend(CANN)
 ggml_add_backend(CUDA)
 ggml_add_backend(HIP)
+ggml_add_backend(IGPU)
 ggml_add_backend(Kompute)
 ggml_add_backend(METAL)
 ggml_add_backend(MUSA)

ggml/src/ggml-backend-reg.cpp

@@ -66,6 +66,10 @@
 #include "ggml-kompute.h"
 #endif
 
+#ifdef GGML_USE_IGPU
+#include "ggml-igpu.h"
+#endif
+
 // disable C++17 deprecation warning for std::codecvt_utf8
 #if defined(__clang__)
 #    pragma clang diagnostic push
@@ -153,6 +157,9 @@ struct ggml_backend_registry {
     std::vector<ggml_backend_dev_t> devices;
 
     ggml_backend_registry() {
+#ifdef GGML_USE_IGPU
+        register_backend(ggml_backend_igpu_reg());
+#endif
 #ifdef GGML_USE_CUDA
         register_backend(ggml_backend_cuda_reg());
 #endif

ggml/src/ggml-igpu/CMakeLists.txt

@@ -0,0 +1,61 @@
+if (WIN32)
+    message(FATAL_ERROR "Not teste on Windows OS")
+endif()
+
+# pas encore bon dans 
+# => /usr/share/cmake/Modules/CMakeDetermineHIPCompiler.cmake
+# => /usr/lib64/cmake/hip-lang/hip-lang-config.cmake vs /usr/local/lib64/cmake/hip-lang/hip-lang-config.cmake
+#   hipcc                         -v -print-targets | grep "Found HIP installation"  => mauvais?
+#   /usr/lib64/llvm18/bin/clang++ -v -print-targets | grep "Found HIP installation"  => mauvais !!!
+#   clang-18                      -v -print-targets | grep "Found HIP installation"  => OK
+#   hipconfig --rocmpath   => OK  => on va le forcer
+
+# forcer la dernier methode...
+if(NOT CMAKE_HIP_COMPILER_ROCM_ROOT)
+  execute_process(
+    COMMAND hipconfig --rocmpath
+    OUTPUT_VARIABLE _CMAKE_HIPCONFIG_ROCMPATH
+    RESULT_VARIABLE _CMAKE_HIPCONFIG_RESULT
+    )
+  if(_CMAKE_HIPCONFIG_RESULT EQUAL 0 AND EXISTS "${_CMAKE_HIPCONFIG_ROCMPATH}")
+    set(CMAKE_HIP_COMPILER_ROCM_ROOT "${_CMAKE_HIPCONFIG_ROCMPATH}")
+  endif()
+endif()
+
+# Forward AMDGPU_TARGETS to CMAKE_HIP_ARCHITECTURES.
+if (AMDGPU_TARGETS AND NOT CMAKE_HIP_ARCHITECTURES)
+    set(CMAKE_HIP_ARCHITECTURES ${AMDGPU_TARGETS})
+endif()
+cmake_minimum_required(VERSION 3.21)
+
+enable_language(HIP)
+
+find_package(hip     REQUIRED)
+
+# pas testé d'autre version...
+if (${hip_VERSION} VERSION_LESS 6.2)
+    message(FATAL_ERROR "At least ROCM/HIP V6.2 is required")
+endif()
+
+message(STATUS "HIP found")
+
+set(TARGET_NAME ggml-igpu)
+
+file(GLOB   GGML_SOURCES_ROCM "ggml-igpu.cpp")
+#file(GLOB   SRCS "*.hip")
+#list(APPEND GGML_SOURCES_ROCM ${SRCS})
+
+ggml_add_backend_library(${TARGET_NAME}
+                         ../../include/ggml-igpu.h
+                         ${GGML_SOURCES_ROCM}
+                        )
+
+add_compile_definitions(GGML_USE_IGPU)
+
+set_source_files_properties(${GGML_SOURCES_ROCM} PROPERTIES LANGUAGE HIP)
+
+if (GGML_STATIC)
+    message(FATAL_ERROR "Static linking not supported for HIP/ROCm")
+endif()
+
+target_link_libraries(${TARGET_NAME} PRIVATE ggml-base hip::host)