backup

Author: airMeng

ggml-sycl.cpp

@@ -36,8 +36,6 @@ GGML_API GGML_CALL int ggml_backend_sycl_get_device_index(int device_id);
 // TODO: these are temporary
 //       ref: https://github.com/ggerganov/llama.cpp/pull/6022#issuecomment-1992615670
 GGML_API GGML_CALL int ggml_backend_sycl_get_device_id(int device_index);
-GGML_API GGML_CALL void ggml_backend_sycl_set_single_device_mode(int main_gpu_id);
-GGML_API GGML_CALL void ggml_backend_sycl_set_mul_device_mode();
 
 // SYCL doesn't support registering host memory, keep here for reference
 // GGML_API GGML_CALL bool ggml_backend_sycl_register_host_buffer(void * buffer, size_t size);

ggml-sycl.h

@@ -19,6 +19,5 @@
 #include "dmmv.hpp"
 #include "mmq.hpp"
 #include "mmvq.hpp"
-#include "pool.hpp"
 
 #endif // GGML_SYCL_BACKEND_HPP

ggml-sycl/backend.hpp

@@ -20,122 +20,6 @@ int get_current_device_id() {
   return dpct::dev_mgr::instance().current_device_id();
 }
 
-void log_ggml_var_device(
-    const char* name,
-    float* src,
-    size_t total_elements,
-    bool src_on_device) {
-  if (!g_ggml_sycl_debug)
-    return;
-  if (!src) {
-    printf("GGML Tensor:%s skip to save for NULL pointer\n", name);
-    return;
-  }
-  char filename[1024];
-  sprintf(filename, "%s.txt", name);
-  printf("GGML Tensor:%s save to %s\n", name, filename);
-
-  size_t total_size = total_elements * sizeof(float);
-  float* local_buf = NULL;
-  if (src_on_device) {
-    local_buf = (float*)ggml_sycl_host_malloc(total_size);
-    ggml_sycl_set_device(g_main_device);
-    dpct::queue_ptr main_stream = g_syclStreams[g_main_device][0];
-    main_stream->memcpy(local_buf, src, total_size).wait();
-  } else {
-    local_buf = (float*)src;
-  }
-
-  std::ofstream logfile;
-  logfile.open(filename);
-  for (size_t i = 0; i < total_elements; i++) {
-    logfile << local_buf[i] << " ";
-    if ((i + 1) % 20 == 0)
-      logfile << std::endl;
-  }
-  logfile << std::endl;
-  logfile.close();
-
-  if (src_on_device)
-    ggml_sycl_host_free(local_buf);
-}
-
-void log_ggml_var_device_fp16(
-    const char* name,
-    sycl::half* src,
-    size_t total_elements,
-    bool src_on_device) {
-  if (!g_ggml_sycl_debug)
-    return;
-  if (!src) {
-    printf("GGML Tensor:%s skip to save for NULL pointer\n", name);
-    return;
-  }
-  char filename[1024];
-  sprintf(filename, "%s.txt", name);
-  printf("GGML Tensor:%s save to %s\n", name, filename);
-
-  size_t total_size = total_elements * sizeof(sycl::half);
-  sycl::half* local_buf = NULL;
-  if (src_on_device) {
-    local_buf = (sycl::half*)ggml_sycl_host_malloc(total_size);
-    ggml_sycl_set_device(g_main_device);
-    dpct::queue_ptr main_stream = g_syclStreams[g_main_device][0];
-    main_stream->memcpy(local_buf, src, total_size).wait();
-  } else {
-    local_buf = (sycl::half*)src;
-  }
-
-  std::ofstream logfile;
-  logfile.open(filename);
-  for (size_t i = 0; i < total_elements; i++) {
-    logfile << local_buf[i] << " ";
-    if ((i + 1) % 20 == 0)
-      logfile << std::endl;
-  }
-  logfile << std::endl;
-  logfile.close();
-
-  if (src_on_device)
-    ggml_sycl_host_free(local_buf);
-}
-
-void print_ggml_tensor(const char* name, struct ggml_tensor* src) {
-  if (!g_ggml_sycl_debug)
-    return;
-  if (!src) {
-    printf("GGML Tensor:%s skip to save for NULL pointer\n", name);
-    return;
-  }
-
-  size_t total_elements = ggml_nelements(src);
-
-  const bool src_on_device = src->backend == GGML_BACKEND_TYPE_GPU ||
-      src->backend == GGML_BACKEND_TYPE_GPU_SPLIT;
-  float* src_data = NULL;
-  if (src_on_device) {
-    ggml_tensor_extra_gpu* src_extra = (ggml_tensor_extra_gpu*)src->extra;
-    src_data = (float*)src_extra->data_device[g_main_device];
-  } else {
-    src_data = (float*)src->data;
-  }
-
-  log_ggml_var_device(name, src_data, total_elements, src_on_device);
-}
-
-void log_tensor_with_cnt(
-    const char* name,
-    struct ggml_tensor* src,
-    int stop_cnt) {
-  stop_cnt = 4;
-  if (log_file_name_idx >= stop_cnt)
-    return;
-  char filename[1280];
-  sprintf(filename, "%s_%07d", name, log_file_name_idx);
-  log_file_name_idx++;
-  print_ggml_tensor(filename, src);
-}
-
 void* ggml_sycl_host_malloc(size_t size) try {
   if (getenv("GGML_SYCL_NO_PINNED") != nullptr) {
     return nullptr;

ggml-sycl/common.cpp

@@ -78,6 +78,8 @@ static int g_work_group_size = 0;
 #define GGML_SYCL_MMV_Y 1
 #endif
 
+typedef sycl::queue *queue_ptr;
+
 enum ggml_sycl_backend_gpu_mode {
   SYCL_UNSET_GPU_MODE = -1,
   SYCL_SINGLE_GPU_MODE = 0,
@@ -182,17 +184,6 @@ static_assert(
 #endif // GGML_SYCL_PEER_MAX_BATCH_SIZE
 
 #define MUL_MAT_SRC1_COL_STRIDE 128
-#define MAX_STREAMS 8
-#define SYCL_MAX_DEVICES 48
-
-static dpct::queue_ptr g_syclStreams[SYCL_MAX_DEVICES][MAX_STREAMS] = {{0}};
-
-struct ggml_tensor_extra_gpu {
-  void* data_device[SYCL_MAX_DEVICES]; // 1 pointer for each device for split
-                                       // tensors
-  dpct::event_ptr events[SYCL_MAX_DEVICES]
-                        [MAX_STREAMS]; // events for synchronizing multiple GPUs
-};
 
 class sycl_gpu_mgr {
  public:
@@ -320,7 +311,7 @@ class sycl_gpu_mgr {
   }
 };
 
-static sycl_gpu_mgr* g_sycl_gpu_mgr = NULL;
+static sycl_gpu_mgr* g_sycl_gpu_mgr = new sycl_gpu_mgr(0);
 static int g_device_count = -1;
 static int g_all_sycl_device_count = -1;
 static int g_main_device = -1;
@@ -329,31 +320,15 @@ static bool g_ggml_backend_sycl_buffer_type_initialized = false;
 
 static std::array<float, SYCL_MAX_DEVICES> g_default_tensor_split = {};
 
-static float g_tensor_split[SYCL_MAX_DEVICES] = {0};
+static float g_tensor_split[GGML_SYCL_MAX_DEVICES] = {0};
 
 static ggml_sycl_backend_gpu_mode g_ggml_sycl_backend_gpu_mode =
     SYCL_UNSET_GPU_MODE;
 
-struct sycl_device_capabilities {
-  int cc; // compute capability
-  bool vmm; // virtual memory support
-  size_t vmm_granularity; // granularity of virtual memory
-  int device_id;
-};
-
-static sycl_device_capabilities g_device_caps[SYCL_MAX_DEVICES] = {
-    {0, false, 0, -1}};
-
-struct sycl_device_id2index {
-  int index;
-};
-
 static void* g_scratch_buffer = nullptr;
 static size_t g_scratch_size = 0; // disabled by default
 static size_t g_scratch_offset = 0;
 
-static dpct::queue_ptr g_sycl_handles[SYCL_MAX_DEVICES] = {nullptr};
-
 int get_main_device();
 
 [[noreturn]] static inline void bad_arch(const sycl::stream& stream_ct1) {
@@ -427,25 +402,151 @@ inline dpct::err0 ggml_sycl_set_device(const int device) try {
   std::exit(1);
 }
 
-void log_ggml_var_device(
-    const char* name,
-    float* src,
-    size_t total_elements,
-    bool src_on_device);
-
-void log_ggml_var_device_fp16(
-    const char* name,
-    sycl::half* src,
-    size_t total_elements,
-    bool src_on_device);
-
-// todo: debug for crash in some case
-void print_ggml_tensor(const char* name, struct ggml_tensor* src);
-
-static int log_file_name_idx = 0;
-void log_tensor_with_cnt(
-    const char* name,
-    struct ggml_tensor* src,
-    int stop_cnt);
+//////////////////////
+
+struct ggml_sycl_device_info {
+    int device_count;
+
+    struct sycl_device_info {
+        int     cc;                 // compute capability
+        // int     nsm;                // number of streaming multiprocessors
+        // size_t  smpb;               // max. shared memory per block
+        bool    vmm;                // virtual memory support
+        size_t  total_vram;
+    };
+
+    sycl_device_info devices[GGML_SYCL_MAX_DEVICES] = {};
+
+    std::array<float, GGML_SYCL_MAX_DEVICES> default_tensor_split = {};
+};
+
+const ggml_sycl_device_info & ggml_sycl_info();
+
+struct ggml_sycl_pool {
+    virtual ~ggml_sycl_pool() = default;
+
+    virtual void * alloc(size_t size, size_t * actual_size) = 0;
+    virtual void free(void * ptr, size_t size) = 0;
+};
+
+template<typename T>
+struct ggml_sycl_pool_alloc {
+    ggml_sycl_pool * pool = nullptr;
+    T * ptr = nullptr;
+    size_t actual_size = 0;
+
+    explicit ggml_sycl_pool_alloc(ggml_sycl_pool & pool) : pool(&pool) {
+    }
+
+    ggml_sycl_pool_alloc(ggml_sycl_pool & pool, size_t size) : pool(&pool) {
+        alloc(size);
+    }
+
+    ~ggml_sycl_pool_alloc() {
+        if (ptr != nullptr) {
+            pool->free(ptr, actual_size);
+        }
+    }
+
+    // size is in number of elements
+    T * alloc(size_t size) {
+        GGML_ASSERT(pool != nullptr);      
+        GGML_ASSERT(ptr == nullptr);
+        ptr = (T *) pool->alloc(size * sizeof(T), &this->actual_size);
+        return ptr;
+    }
+
+    T * alloc(ggml_sycl_pool & pool, size_t size) {
+        this->pool = &pool;
+        return alloc(size);
+    }
+
+    T * get() {
+        return ptr;
+    }
+
+    ggml_sycl_pool_alloc() = default;
+    ggml_sycl_pool_alloc(const ggml_sycl_pool_alloc &) = delete;
+    ggml_sycl_pool_alloc(ggml_sycl_pool_alloc &&) = delete;
+    ggml_sycl_pool_alloc& operator=(const ggml_sycl_pool_alloc &) = delete;
+    ggml_sycl_pool_alloc& operator=(ggml_sycl_pool_alloc &&) = delete;
+};
+
+// backend interface
+
+struct ggml_tensor_extra_gpu {
+  void* data_device[GGML_SYCL_MAX_DEVICES]; // 1 pointer for each device for split
+                                       // tensors
+  dpct::event_ptr events[GGML_SYCL_MAX_DEVICES]
+                        [GGML_SYCL_MAX_STREAMS]; // events for synchronizing multiple GPUs
+};
+
+struct ggml_backend_sycl_context {
+    int device;
+    std::string name;
+
+    queue_ptr qptrs[GGML_SYCL_MAX_DEVICES][GGML_SYCL_MAX_STREAMS] = { { nullptr } };
+    static sycl::handler * sycl_handles[GGML_SYCL_MAX_DEVICES] = {nullptr};
+
+    explicit ggml_backend_sycl_context(int device) :
+        device(device),
+        name(GGML_SYCL_NAME + std::to_string(device)) {
+    }
+
+    ~ggml_backend_sycl_context() {
+        for (int i = 0; i < GGML_SYCL_MAX_DEVICES; ++i) {
+            for (int j = 0; j < GGML_SYCL_MAX_STREAMS; ++j) {
+                if (streams[i][j] != nullptr) {
+                    SYCL_CHECK(free(streams[i][j]));
+                }
+            }
+            if (cublas_handles[i] != nullptr) {
+                SYCL_CHECK(free(sycl_handles[i]));
+            }
+        }
+    }
+
+    queue_ptr stream(int device, int stream) {
+        if (qptrs[device][stream] == nullptr) {
+            SYCL_CHECK(dpct::get_current_device().create_queue(
+                        g_sycl_gpu_mgr->get_co_ctx(), dpct::get_current_device())));
+        }
+        return qptrs[device][stream];
+    }
+
+    cudaStream_t stream() {
+        return stream(device, 0);
+    }
+
+    cublasHandle_t sycl_handle(int device) {
+        if (sycl_handles[device] == nullptr) {
+            const dpct::queue_ptr stream = streams[device][0];
+            // create sycl handle
+            SYCL_CHECK(CHECK_TRY_ERROR(sycl_handles[device] = stream));
+        }
+        return sycl_handles[device];
+    }
+
+    cublasHandle_t sycl_handle() {
+        return sycl_handle(device);
+    }
+
+    // pool
+    std::unique_ptr<ggml_sycl_pool> pools[GGML_SYCL_MAX_DEVICES];
+
+    static std::unique_ptr<ggml_sycl_pool> new_pool_for_device(queue_ptr qptr, int device);
+
+    ggml_sycl_pool & pool(int device) {
+        if (pools[device] == nullptr) {
+            pools[device] = new_pool_for_device(qptrs[device][0], device);
+        }
+        return *pools[device];
+    }
+
+    ggml_sycl_pool & pool() {
+        return pool(device);
+    }
+};
+
 
 #endif // GGML_SYCL_COMMON_HPP

ggml-sycl/common.hpp

@@ -1,6 +1,6 @@
 #include "convert.hpp"
 #include "dequantize.hpp"
-
+#include "presets.hpp"
 
 template <int qk, int qr, dequantize_kernel_t dequantize_kernel, typename dst_t>
 static void dequantize_block(const void * __restrict__ vx, dst_t * __restrict__ y, const int k,