Mirrors/llama.cpp
1
0
Fork 0
mirror of https://github.com/ggerganov/llama.cpp.git synced 2025-01-13 13:52:22 +01:00
Code Issues Packages Projects Releases Wiki Activity

Replace powf with sycl::pow in ggml-sycl.cpp

Browse Source 
Signed-off-by: Joe Todd <joe.todd@codeplay.com>
This commit is contained in:
Joe Todd 2024-06-14 13:14:33 +01:00
parent abd7c7b8c2
commit ded54b5d9b
1 changed files with 6 additions and 6 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

12
ggml-sycl.cpp
View File

@ -8877,7 +8877,7 @@ static void rope_norm(
const int i = row*ne0 + i0;
const int i2 = row/p_delta_rows;
const float theta_base = pos[i2]*powf(theta_scale, i0/2.0f);
const float theta_base = pos[i2]*sycl::pow(theta_scale, i0/2.0f);
const float freq_factor = has_ff ? freq_factors[i0 / 2] : 1.0f;
float cos_theta, sin_theta;
@ -8919,7 +8919,7 @@ static void rope_neox(const T *x, T *dst, int ne0, int n_dims,
const int i = row*ne0 + i0/2;
const int i2 = row/p_delta_rows;
const float theta_base = pos[i2]*powf(theta_scale, i0/2.0f);
const float theta_base = pos[i2]*sycl::pow(theta_scale, i0/2.0f);
const float freq_factor = has_ff ? freq_factors[i0/2] : 1.0f;
float cos_theta, sin_theta;
@ -12388,7 +12388,7 @@ static void rope_norm_sycl(const T *x, T *dst, int ne0, int n_dims, int nr,
const int n_blocks_x = (ne0 + 2*SYCL_ROPE_BLOCK_SIZE - 1) / (2*SYCL_ROPE_BLOCK_SIZE);
const sycl::range<3> block_nums(1, n_blocks_x, nr);
const float theta_scale = powf(freq_base, -2.0f/n_dims);
const float theta_scale = sycl::pow(freq_base, -2.0f/n_dims);
if (freq_factors == nullptr) {
/*
@ -12436,7 +12436,7 @@ static void rope_neox_sycl(const T *x, T *dst, int ne0, int n_dims, int nr,
const int n_blocks_x = (ne0 + 2*SYCL_ROPE_BLOCK_SIZE - 1) / (2*SYCL_ROPE_BLOCK_SIZE);
const sycl::range<3> block_nums(1, n_blocks_x, nr);
const float theta_scale = powf(freq_base, -2.0f/n_dims);
const float theta_scale = sycl::pow(freq_base, -2.0f/n_dims);
dpct::has_capability_or_fail(stream->get_device(),
{sycl::aspect::fp16});
@ -12575,8 +12575,8 @@ static void soft_max_f32_sycl(const float * x, const float * mask,
const uint32_t n_head_kv = nrows_x/nrows_y;
const uint32_t n_head_log2 = 1u << (uint32_t) floorf(log2f((float) n_head_kv));
const float m0 = powf(2.0f, -(max_bias ) / n_head_log2);
const float m1 = powf(2.0f, -(max_bias / 2.0f) / n_head_log2);
const float m0 = sycl::pow(2.0f, -(max_bias ) / n_head_log2);
const float m1 = sycl::pow(2.0f, -(max_bias / 2.0f) / n_head_log2);
const size_t local_mem_size = stream->get_device().get_info<sycl::info::device::local_mem_size>();
if (n_local_scratch*sizeof(float) < local_mem_size) {