@@ -12438,26 +12438,26 @@ static void rope_neox_sycl(const T *x, T *dst, int ne0, int n_dims, int nr,
1243812438
1243912439 const float theta_scale = powf(freq_base, -2.0f/n_dims);
1244012440
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) {
1244712447 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,
1244912449 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) {
1245612456 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,
1245812458 corr_dims, theta_scale, freq_factors,
12459- item_ct1);
12460- });
12459+ item_ct1);
12460+ });
1246112461 }
1246212462}
1246312463
@@ -14010,8 +14010,8 @@ inline void ggml_sycl_op_rope(const ggml_tensor *src0, const ggml_tensor *src1,
1401014010 const int32_t * pos = (const int32_t *) src1_dd;
1401114011
1401214012 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;
1401514015 }
1401614016
1401714017 rope_corr_dims corr_dims;
0 commit comments