Fix bugs in host_mem pool implementation.

First fully working implementation in benchmark and llama-cli. Still
issues in backend tests.

Signed-off-by: nscipione <[email protected]>

Author: s-Nick

ggml/src/ggml-sycl/common.cpp

@@ -29,7 +29,7 @@ void* ggml_sycl_host_malloc(size_t size) try {
   dpct::err0 err = CHECK_TRY_ERROR(
       ptr = (void*)sycl::malloc_host(size, dpct::get_in_order_queue()));
 
-    printf("Luigi\n");
+    //printf("Luigi\n");
   if (err != 0) {
     // clear the error
     GGML_LOG_ERROR("WARNING: failed to allocate %.2f MB of pinned memory: %s\n", size / 1024.0 / 1024.0,    "syclGetErrorString is not supported");

ggml/src/ggml-sycl/dpct/helper.hpp

@@ -18,6 +18,7 @@
 #include <syclcompat/math.hpp>
 #include <oneapi/mkl.hpp>
 #include <map>
+#include <cassert>
 
 #include "ggml-sycl.h"
 #include "ggml.h"
@@ -88,6 +89,16 @@ inline std::string get_device_backend_and_type(const sycl::device &device) {
     return device_type.str();
 }
 
+template<typename Ts>
+struct matrix_info_t
+{
+    oneapi::mkl::transpose transpose_info[2];
+    Ts value_info[2];
+    std::int64_t size_info[3];
+    std::int64_t ld_info[3];
+    std::int64_t groupsize_info;
+};
+
 namespace dpct
 {
     typedef sycl::queue *queue_ptr;
@@ -1737,27 +1748,16 @@ namespace dpct
                                     oneapi::mkl::transpose b_trans, int m, int n, int k,
                                     const void *alpha, const void **a, int lda,
                                     const void **b, int ldb, const void *beta, void **c,
-                                    int ldc, int batch_size)
+                                    int ldc, int batch_size, matrix_info_t<double>* matrix_info)
         {
-            struct matrix_info_t
-            {
-                oneapi::mkl::transpose transpose_info[2];
-                Ts value_info[2];
-                std::int64_t size_info[3];
-                std::int64_t ld_info[3];
-                std::int64_t groupsize_info;
-            };
 
             Ts alpha_value = dpct::get_value(reinterpret_cast<const Ts *>(alpha), q);
             Ts beta_value = dpct::get_value(reinterpret_cast<const Ts *>(beta), q);
 
-            //ggml_backend_sycl_host_buffer_type()->alloc_buffer;
-            auto tmp = ggml_backend_sycl_reg();
-            std::cout << "this is WARIO " << tmp->iface.get_name(tmp) << '\n';
+            //::matrix_info_t<Ts> *matrix_info =
+                //(::matrix_info_t<Ts> *)std::malloc(sizeof(matrix_info_t<Ts>));
+            //printf("test pointer %p alpha_value %f before\n", matrix_info, alpha_value);
 
-
-            matrix_info_t *matrix_info =
-                (matrix_info_t *)std::malloc(sizeof(matrix_info_t));
             matrix_info->transpose_info[0] = a_trans;
             matrix_info->transpose_info[1] = b_trans;
             matrix_info->value_info[0] = alpha_value;
@@ -1770,13 +1770,15 @@ namespace dpct
             matrix_info->ld_info[2] = ldc;
             matrix_info->groupsize_info = batch_size;
 
+            //printf("test pointer %p alpha_value %f\n", matrix_info, matrix_info->value_info[0]);;
+
 #ifdef GGML_SYCL_NVIDIA
             sycl::event e = oneapi::mkl::blas::column_major::gemm_batch(
                 oneapi::mkl::backend_selector<oneapi::mkl::backend::cublas>{ q }, matrix_info->transpose_info,
                 matrix_info->transpose_info + 1, matrix_info->size_info, matrix_info->size_info + 1,
-                matrix_info->size_info + 2, matrix_info->value_info, reinterpret_cast<const Ta **>(a),
+                matrix_info->size_info + 2, reinterpret_cast<Ts*>(matrix_info->value_info), reinterpret_cast<const Ta **>(a),
                 matrix_info->ld_info, reinterpret_cast<const Tb **>(b), matrix_info->ld_info + 1,
-                matrix_info->value_info + 1, reinterpret_cast<Tc **>(c), matrix_info->ld_info + 2, 1,
+                reinterpret_cast<Ts*>(matrix_info->value_info + 1), reinterpret_cast<Tc **>(c), matrix_info->ld_info + 2, 1,
                 &(matrix_info->groupsize_info));
 #else
             sycl::event e = oneapi::mkl::blas::column_major::gemm_batch(
@@ -1786,11 +1788,14 @@ namespace dpct
                 matrix_info->ld_info + 1, matrix_info->value_info + 1, reinterpret_cast<Tc **>(c),
                 matrix_info->ld_info + 2, 1, &(matrix_info->groupsize_info));
 #endif
+        //printf("gemm_launched\n");
 
+    /*
             q.submit([&](sycl::handler &cgh)
                      {
     cgh.depends_on(e);
     cgh.host_task([=] { std::free(matrix_info); }); });
+        */
         }
 
         template <class Ta, class Tb, class Tc, class Ts>
@@ -2439,7 +2444,8 @@ namespace dpct
                            library_data_t a_type, int lda, const void *b[],
                            library_data_t b_type, int ldb, const void *beta,
                            void *c[], library_data_t c_type, int ldc,
-                           int batch_size, library_data_t scaling_type)
+                           int batch_size, library_data_t scaling_type,
+                           matrix_info_t<double>* matrix_info)
     {
         if (scaling_type == library_data_t::real_float &&
             c_type == library_data_t::complex_float)
@@ -2451,7 +2457,6 @@ namespace dpct
         {
             scaling_type = library_data_t::complex_double;
         }
-
         std::uint64_t key =
             detail::get_type_combination_id(a_type, b_type, c_type, scaling_type);
         switch (key)
@@ -2462,7 +2467,7 @@ namespace dpct
         {
             detail::gemm_batch_impl<float, float, float, float>(
                 q, a_trans, b_trans, m, n, k, alpha, a, lda, b, ldb, beta, c, ldc,
-                batch_size);
+                batch_size, matrix_info);
             break;
         }
         case detail::get_type_combination_id(
@@ -2471,17 +2476,18 @@ namespace dpct
         {
             detail::gemm_batch_impl<double, double, double, double>(
                 q, a_trans, b_trans, m, n, k, alpha, a, lda, b, ldb, beta, c, ldc,
-                batch_size);
+                batch_size, matrix_info);
             break;
         }
+/*
         case detail::get_type_combination_id(
             library_data_t::complex_float, library_data_t::complex_float,
             library_data_t::complex_float, library_data_t::complex_float):
         {
             detail::gemm_batch_impl<std::complex<float>, std::complex<float>,
                                     std::complex<float>, std::complex<float>>(
                 q, a_trans, b_trans, m, n, k, alpha, a, lda, b, ldb, beta, c, ldc,
-                batch_size);
+                batch_size, matrix_info);
             break;
         }
         case detail::get_type_combination_id(
@@ -2491,17 +2497,18 @@ namespace dpct
             detail::gemm_batch_impl<std::complex<double>, std::complex<double>,
                                     std::complex<double>, std::complex<double>>(
                 q, a_trans, b_trans, m, n, k, alpha, a, lda, b, ldb, beta, c, ldc,
-                batch_size);
+                batch_size, matrix_info);
             break;
         }
+*/
         case detail::get_type_combination_id(
             library_data_t::real_half, library_data_t::real_half,
             library_data_t::real_half, library_data_t::real_half):
         {
             detail::gemm_batch_impl<sycl::half, sycl::half, sycl::half,
                                     sycl::half>(q, a_trans, b_trans, m, n, k, alpha,
                                                 a, lda, b, ldb, beta, c, ldc,
-                                                batch_size);
+                                                batch_size, matrix_info);
             break;
         }
 #ifdef __INTEL_MKL__
@@ -2512,7 +2519,7 @@ namespace dpct
             detail::gemm_batch_impl<oneapi::mkl::bfloat16, oneapi::mkl::bfloat16,
                                     oneapi::mkl::bfloat16, float>(
                 q, a_trans, b_trans, m, n, k, alpha, a, lda, b, ldb, beta, c, ldc,
-                batch_size);
+                batch_size, matrix_info);
             break;
         }
         case detail::get_type_combination_id(
@@ -2521,7 +2528,7 @@ namespace dpct
         {
             detail::gemm_batch_impl<oneapi::mkl::bfloat16, oneapi::mkl::bfloat16, float,
                                     float>(q, a_trans, b_trans, m, n, k, alpha, a, lda,
-                                           b, ldb, beta, c, ldc, batch_size);
+                                           b, ldb, beta, c, ldc, batch_size, matrix_info);
             break;
         }
 #endif
@@ -2536,7 +2543,7 @@ namespace dpct
             detail::gemm_batch_impl<std::int8_t, std::int8_t, std::int32_t,
                                     float>(q, a_trans, b_trans, m, n, k, &alpha_float,
                                            a, lda, b, ldb, &beta_float, c, ldc,
-                                           batch_size);
+                                           batch_size, matrix_info);
             break;
         }
         case detail::get_type_combination_id(
@@ -2545,7 +2552,7 @@ namespace dpct
         {
             detail::gemm_batch_impl<std::int8_t, std::int8_t, float, float>(
                 q, a_trans, b_trans, m, n, k, alpha, a, lda, b, ldb, beta, c, ldc,
-                batch_size);
+                batch_size, matrix_info);
             break;
         }
         case detail::get_type_combination_id(
@@ -2554,7 +2561,7 @@ namespace dpct
         {
             detail::gemm_batch_impl<sycl::half, sycl::half, float, float>(
                 q, a_trans, b_trans, m, n, k, alpha, a, lda, b, ldb, beta, c, ldc,
-                batch_size);
+                batch_size, matrix_info);
             break;
         }
         case detail::get_type_combination_id(
@@ -2569,7 +2576,7 @@ namespace dpct
             sycl::half beta_half(beta_value);
             detail::gemm_batch_impl<sycl::half, sycl::half, sycl::half, sycl::half>(
                 q, a_trans, b_trans, m, n, k, &alpha_half, a, lda, b, ldb, &beta_half, c, ldc,
-                batch_size);
+                batch_size, matrix_info);
             break;
         }
         default: