Caching device_info in device_ext to avoid repetitive queries

ggerganov · Jul 4, 2024 · 8055548 · 8055548
1 parent 807b0c4
commit 8055548
Show file tree

Hide file tree

Showing 4 changed files with 23 additions and 21 deletions.
diff --git a/ggml/src/ggml-sycl.cpp b/ggml/src/ggml-sycl.cpp
@@ -49,7 +49,6 @@ bool   ggml_backend_is_sycl(ggml_backend_t backend);
 int    ggml_backend_sycl_get_device(ggml_backend_t backend);
 static bool ggml_backend_buffer_is_sycl_split(ggml_backend_buffer_t buffer);
 static inline int get_sycl_env(const char *env_name, int default_val);
-static inline int get_work_group_size(const sycl::device& device);
 
 void dev2dev_memcpy(sycl::queue &q_dst, sycl::queue &q_src, void *ptr_dst,
                     const void *ptr_src, size_t size) {
@@ -1914,7 +1913,7 @@ static void soft_max_f32_sycl(const float * x, const float * mask,
                               const int nrows_y, const float scale, const float max_bias,
                               queue_ptr stream) {
     int nth = WARP_SIZE;
-    int max_block_size = get_work_group_size(stream->get_device());
+    int max_block_size = dpct::dev_mgr::instance().get_work_group_size(stream->get_device());
     while (nth < ncols_x && nth < max_block_size) nth *= 2;
     if (nth>max_block_size) nth = max_block_size;
 

diff --git a/ggml/src/ggml-sycl/common.hpp b/ggml/src/ggml-sycl/common.hpp
@@ -295,15 +295,6 @@ struct ggml_backend_sycl_context {
     }
 };
 
-// common host functions
-
-static inline int get_work_group_size(const sycl::device& device) {
-    dpct::device_info prop;
-    dpct::get_device_info(prop, device);
-    return prop.get_max_work_group_size();
-}
-
-
 // common device functions
 
 static __dpct_inline__ float warp_reduce_sum(float x,

diff --git a/ggml/src/ggml-sycl/dpct/helper.hpp b/ggml/src/ggml-sycl/dpct/helper.hpp
@@ -672,13 +672,16 @@ namespace dpct
       }
 
       void get_device_info(device_info &out) const {
-        dpct::get_device_info(out, *this);
+        out = this->get_device_info();
       }
 
-      device_info get_device_info() const {
-        device_info prop;
-        dpct::get_device_info(prop, *this);
-        return prop;
+      const device_info& get_device_info() const {
+        std::lock_guard<std::mutex> lock(m_mutex);
+        if (!_dev_info) {
+            _dev_info = device_info{};
+            dpct::get_device_info(*_dev_info, *this);
+        }
+        return *_dev_info;
       }
 
       void reset() {
@@ -801,6 +804,7 @@ namespace dpct
       sycl::queue _saved_queue;
       std::vector<sycl::queue> _queues;
       mutable mutex_type m_mutex;
+      mutable std::optional<device_info> _dev_info;
     };
 
 
@@ -852,7 +856,7 @@ namespace dpct
         }
         unsigned int device_count() { return _devs.size(); }
 
-        unsigned int get_device_id(const sycl::device &dev)
+        unsigned int get_device_id(const sycl::device &dev) const
         {
             unsigned int id = 0;
             for (auto dev_item : _devs)
@@ -882,6 +886,15 @@ namespace dpct
             static dev_mgr d_m;
             return d_m;
         }
+
+        int get_work_group_size(unsigned int id) const {
+            return get_device(id).get_max_work_group_size();
+        }
+
+        int get_work_group_size(const sycl::device &dev) const {
+            return get_work_group_size(get_device_id(dev));
+        }
+
         dev_mgr(const dev_mgr &) = delete;
         dev_mgr &operator=(const dev_mgr &) = delete;
         dev_mgr(dev_mgr &&) = delete;
@@ -2623,7 +2636,6 @@ namespace dpct
                 beta, c, ldc, stride_c, batch_size);
             break;
         }
-#endif
         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_float):

diff --git a/ggml/src/ggml-sycl/norm.cpp b/ggml/src/ggml-sycl/norm.cpp
@@ -197,7 +197,7 @@ static void norm_f32_sycl(const float* x, float* dst, const int ncols,
             });
     }
     else {
-        const int work_group_size = get_work_group_size(stream->get_device());
+        const int work_group_size = dpct::dev_mgr::instance().get_work_group_size(stream->get_device());
         const sycl::range<3> block_dims(1, 1, work_group_size);
         /*
         DPCT1049:17: The work-group size passed to the SYCL kernel may exceed
@@ -240,7 +240,7 @@ static void group_norm_f32_sycl(const float* x, float* dst,
             });
     }
     else {
-        const int work_group_size = get_work_group_size(stream->get_device());
+        const int work_group_size = dpct::dev_mgr::instance().get_work_group_size(stream->get_device());
         const sycl::range<3> block_dims(1, 1, work_group_size);
         /*
         DPCT1049:18: The work-group size passed to the SYCL kernel may exceed
@@ -286,7 +286,7 @@ static void rms_norm_f32_sycl(const float* x, float* dst, const int ncols,
             });
     }
     else {
-        const int work_group_size = get_work_group_size(stream->get_device());
+        const int work_group_size = dpct::dev_mgr::instance().get_work_group_size(stream->get_device());
         const sycl::range<3> block_dims(1, 1, work_group_size);
         /*
         DPCT1049:19: The work-group size passed to the SYCL kernel may exceed