vbatts/llama.cpp
1
0
Fork 0
Code Issues Pull requests Projects Releases Packages Wiki Activity Actions

Caching device_info in device_ext to avoid repetitive queries

Browse source
This commit is contained in:
OuadiElfarouki 2024-07-04 16:30:34 +01:00
parent 807b0c49ff
commit 80555483aa
4 changed files with 23 additions and 21 deletions
Download patch file Download diff file Expand all files Collapse all files

3
ggml/src/ggml-sycl.cpp
View file

@ -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;

9
ggml/src/ggml-sycl/common.hpp
View file

@ -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,

26
ggml/src/ggml-sycl/dpct/helper.hpp
View file

@ -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):

6
ggml/src/ggml-sycl/norm.cpp
View file

@ -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