Skip to content

Commit

Permalink
caching device_info in device_ext to avoid repetitive queries
Browse files Browse the repository at this point in the history
  • Loading branch information
@OuadiElfarouki
OuadiElfarouki committed Jul 4, 2024
1 parent 807b0c4 commit 3424fce
Show file tree
Hide file tree
Showing 4 changed files with 25 additions and 52 deletions.
33 changes: 1 addition & 32 deletions ggml/src/ggml-sycl.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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) {
Expand Down Expand Up @@ -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;

Expand Down Expand Up @@ -3729,10 +3728,6 @@ static void ggml_sycl_mul_mat_batched_sycl(ggml_backend_sycl_context & ctx,
SYCL_CHECK(ggml_sycl_set_device(ctx.device));
queue_ptr main_stream = ctx.stream();;

bool no_mixed_dtypes = main_stream->get_backend() == sycl::backend::ext_oneapi_cuda ||
main_stream->get_backend() == sycl::backend::ext_oneapi_hip;


void * src0_ddq = src0->data;
sycl::half *src0_as_f16 = (sycl::half *)src0_ddq;
float * src1_ddf = (float *) src1->data;
Expand All @@ -3750,15 +3745,10 @@ static void ggml_sycl_mul_mat_batched_sycl(ggml_backend_sycl_context & ctx,
sycl::half *src1_f16 = src1->type == GGML_TYPE_F16 ? (sycl::half *)src1_ddf
: src1_f16_alloc.get();

ggml_sycl_pool_alloc<sycl::half> dst_f16(ctx.pool());
char * dst_t;

dpct::library_data_t cu_compute_type = dpct::library_data_t::real_float;
dpct::library_data_t cu_data_type = dpct::library_data_t::real_float;
if (no_mixed_dtypes) {
cu_compute_type = dpct::library_data_t::real_half;
cu_data_type = dpct::library_data_t::real_half;
}

// dst strides
size_t nbd2 = dst->nb[2];
Expand All @@ -3767,26 +3757,10 @@ static void ggml_sycl_mul_mat_batched_sycl(ggml_backend_sycl_context & ctx,
const float alpha_f32 = 1.0f;
const float beta_f32 = 0.0f;

const sycl::half alpha_f16 = 1.0f;
const sycl::half beta_f16 = 0.0f;

const void * alpha = &alpha_f32;
const void * beta = &beta_f32;
if (no_mixed_dtypes) {
alpha = &alpha_f16;
beta = &beta_f16;
}

// TODO: Renable (dst->op_params[0] =! GGML_PREC_DEFAULT) pathway
// when oneMKL open source supports half, half, float, float: datatypes

dst_t = (char *) dst_ddf;
if (no_mixed_dtypes) {
dst_t = (char *) dst_f16.alloc(ne_dst);

nbd2 /= sizeof(float) / sizeof(sycl::half);
nbd3 /= sizeof(float) / sizeof(sycl::half);
}

GGML_ASSERT(ne12 % ne02 == 0);
GGML_ASSERT(ne13 % ne03 == 0);
Expand Down Expand Up @@ -3848,11 +3822,6 @@ static void ggml_sycl_mul_mat_batched_sycl(ggml_backend_sycl_context & ctx,
(void **)(ptrs_dst.get() + 0 * ne23), cu_data_type, ne01, ne23,
cu_compute_type)));
}

if (no_mixed_dtypes) {
const to_fp32_sycl_t to_fp32_sycl = ggml_get_to_fp32_sycl(GGML_TYPE_F16);
to_fp32_sycl(dst_f16.get(), dst_ddf, ne_dst, main_stream);
}
}
catch (sycl::exception const &exc) {
std::cerr << exc.what() << "Exception caught at file:" << __FILE__
Expand Down
9 changes: 0 additions & 9 deletions ggml/src/ggml-sycl/common.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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,
Expand Down
29 changes: 21 additions & 8 deletions ggml/src/ggml-sycl/dpct/helper.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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() {
Expand Down Expand Up @@ -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;
};
Expand Down Expand Up @@ -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)
Expand Down Expand Up @@ -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;
Expand Down Expand Up @@ -2426,6 +2439,7 @@ namespace dpct
b, ldb, beta, c, ldc, batch_size);
break;
}
#endif
case detail::get_type_combination_id(
library_data_t::real_int8, library_data_t::real_int8,
library_data_t::real_int32, library_data_t::real_int32):
Expand Down Expand Up @@ -2458,7 +2472,6 @@ namespace dpct
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):
Expand Down Expand Up @@ -2595,6 +2608,7 @@ namespace dpct
stride_c, batch_size);
break;
}
#endif
case detail::get_type_combination_id(
library_data_t::real_int8, library_data_t::real_int8,
library_data_t::real_int32, library_data_t::real_int32):
Expand Down Expand Up @@ -2623,7 +2637,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):
Expand Down
6 changes: 3 additions & 3 deletions ggml/src/ggml-sycl/norm.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down Expand Up @@ -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
Expand Down Expand Up @@ -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
Expand Down

0 comments on commit 3424fce

Please sign in to comment.