@@ -12438,26 +12438,26 @@ static void rope_neox_sycl(const T *x, T *dst, int ne0, int n_dims, int nr,
12438
12438
12439
12439
const float theta_scale = powf(freq_base, -2.0f/n_dims);
12440
12440
12441
- dpct::has_capability_or_fail(stream->get_device(),
12442
- {sycl::aspect::fp16});
12443
- if (freq_factors == nullptr) {
12444
- stream->parallel_for(
12445
- sycl::nd_range<3>(block_nums * block_dims, block_dims),
12446
- [=](sycl::nd_item<3> item_ct1) {
12441
+ dpct::has_capability_or_fail(stream->get_device(),
12442
+ {sycl::aspect::fp16});
12443
+ if (freq_factors == nullptr) {
12444
+ stream->parallel_for(
12445
+ sycl::nd_range<3>(block_nums * block_dims, block_dims),
12446
+ [=](sycl::nd_item<3> item_ct1) {
12447
12447
rope_neox<T, false>(x, dst, ne0, n_dims, pos, freq_scale,
12448
- p_delta_rows, ext_factor, attn_factor,
12448
+ p_delta_rows, ext_factor, attn_factor,
12449
12449
corr_dims, theta_scale, freq_factors,
12450
- item_ct1);
12451
- });
12452
- } else {
12453
- stream->parallel_for(
12454
- sycl::nd_range<3>(block_nums * block_dims, block_dims),
12455
- [=](sycl::nd_item<3> item_ct1) {
12450
+ item_ct1);
12451
+ });
12452
+ } else {
12453
+ stream->parallel_for(
12454
+ sycl::nd_range<3>(block_nums * block_dims, block_dims),
12455
+ [=](sycl::nd_item<3> item_ct1) {
12456
12456
rope_neox<T, true>(x, dst, ne0, n_dims, pos, freq_scale,
12457
- p_delta_rows, ext_factor, attn_factor,
12457
+ p_delta_rows, ext_factor, attn_factor,
12458
12458
corr_dims, theta_scale, freq_factors,
12459
- item_ct1);
12460
- });
12459
+ item_ct1);
12460
+ });
12461
12461
}
12462
12462
}
12463
12463
@@ -14010,8 +14010,8 @@ inline void ggml_sycl_op_rope(const ggml_tensor *src0, const ggml_tensor *src1,
14010
14010
const int32_t * pos = (const int32_t *) src1_dd;
14011
14011
14012
14012
const float * freq_factors = nullptr;
14013
- if (src2 != nullptr) {
14014
- freq_factors = (const float *) src2->data;
14013
+ if (src2 != nullptr) {
14014
+ freq_factors = (const float *) src2->data;
14015
14015
}
14016
14016
14017
14017
rope_corr_dims corr_dims;
0 commit comments