@@ -8877,7 +8877,7 @@ static void rope_norm(
8877
8877
const int i = row*ne0 + i0;
8878
8878
const int i2 = row/p_delta_rows;
8879
8879
8880
- const float theta_base = pos[i2]*powf (theta_scale, i0/2.0f);
8880
+ const float theta_base = pos[i2]*sycl::pow (theta_scale, i0/2.0f);
8881
8881
const float freq_factor = has_ff ? freq_factors[i0 / 2] : 1.0f;
8882
8882
8883
8883
float cos_theta, sin_theta;
@@ -8919,7 +8919,7 @@ static void rope_neox(const T *x, T *dst, int ne0, int n_dims,
8919
8919
const int i = row*ne0 + i0/2;
8920
8920
const int i2 = row/p_delta_rows;
8921
8921
8922
- const float theta_base = pos[i2]*powf (theta_scale, i0/2.0f);
8922
+ const float theta_base = pos[i2]*sycl::pow (theta_scale, i0/2.0f);
8923
8923
const float freq_factor = has_ff ? freq_factors[i0/2] : 1.0f;
8924
8924
8925
8925
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,
12388
12388
const int n_blocks_x = (ne0 + 2*SYCL_ROPE_BLOCK_SIZE - 1) / (2*SYCL_ROPE_BLOCK_SIZE);
12389
12389
const sycl::range<3> block_nums(1, n_blocks_x, nr);
12390
12390
12391
- const float theta_scale = powf (freq_base, -2.0f/n_dims);
12391
+ const float theta_scale = sycl::pow (freq_base, -2.0f/n_dims);
12392
12392
12393
12393
if (freq_factors == nullptr) {
12394
12394
/*
@@ -12436,7 +12436,7 @@ static void rope_neox_sycl(const T *x, T *dst, int ne0, int n_dims, int nr,
12436
12436
const int n_blocks_x = (ne0 + 2*SYCL_ROPE_BLOCK_SIZE - 1) / (2*SYCL_ROPE_BLOCK_SIZE);
12437
12437
const sycl::range<3> block_nums(1, n_blocks_x, nr);
12438
12438
12439
- const float theta_scale = powf (freq_base, -2.0f/n_dims);
12439
+ const float theta_scale = sycl::pow (freq_base, -2.0f/n_dims);
12440
12440
12441
12441
dpct::has_capability_or_fail(stream->get_device(),
12442
12442
{sycl::aspect::fp16});
@@ -12575,8 +12575,8 @@ static void soft_max_f32_sycl(const float * x, const float * mask,
12575
12575
const uint32_t n_head_kv = nrows_x/nrows_y;
12576
12576
const uint32_t n_head_log2 = 1u << (uint32_t) floorf(log2f((float) n_head_kv));
12577
12577
12578
- const float m0 = powf (2.0f, -(max_bias ) / n_head_log2);
12579
- const float m1 = powf (2.0f, -(max_bias / 2.0f) / n_head_log2);
12578
+ const float m0 = sycl::pow (2.0f, -(max_bias ) / n_head_log2);
12579
+ const float m1 = sycl::pow (2.0f, -(max_bias / 2.0f) / n_head_log2);
12580
12580
12581
12581
const size_t local_mem_size = stream->get_device().get_info<sycl::info::device::local_mem_size>();
12582
12582
if (n_local_scratch*sizeof(float) < local_mem_size) {
0 commit comments