From 80555483aabe7c066755c41c679636fbdf63e27e Mon Sep 17 00:00:00 2001 From: OuadiElfarouki Date: Thu, 4 Jul 2024 16:30:34 +0100 Subject: [PATCH] Caching device_info in device_ext to avoid repetitive queries --- ggml/src/ggml-sycl.cpp | 3 +-- ggml/src/ggml-sycl/common.hpp | 9 --------- ggml/src/ggml-sycl/dpct/helper.hpp | 26 +++++++++++++++++++------- ggml/src/ggml-sycl/norm.cpp | 6 +++--- 4 files changed, 23 insertions(+), 21 deletions(-) diff --git a/ggml/src/ggml-sycl.cpp b/ggml/src/ggml-sycl.cpp index 76bad57e2..43ad9b0e8 100644 --- 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 index 476d847ca..30701dd61 100644 --- 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 index 5e98660dc..db6d9fecd 100644 --- 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 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 _queues; mutable mutex_type m_mutex; + mutable std::optional _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 index a77f7852c..3370ec610 100644 --- 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