@@ -602,11 +602,21 @@ T shift_sub_group_left(sycl::sub_group sg, T input, int delta, int last_item,
602602 return result;
603603#else
604604 if ((1U << sg.get_local_linear_id ()) & member_mask) {
605+ #if defined(__INTEL_LLVM_COMPILER) && __INTEL_LLVM_COMPILER < 20260000
606+ auto inner = sycl::ext::oneapi::experimental::get_tangle_group (sg);
607+ #else
605608 auto inner = sycl::ext::oneapi::experimental::entangle (sg);
609+ #endif
606610 sycl::group_barrier (inner);
607611 }
612+ #if defined(__INTEL_LLVM_COMPILER) && __INTEL_LLVM_COMPILER < 20260000
613+ auto partition =
614+ sycl::ext::oneapi::experimental::get_fixed_size_group<LogicSubGroupSize>(
615+ sg);
616+ #else
608617 auto partition =
609618 sycl::ext::oneapi::experimental::chunked_partition<LogicSubGroupSize>(sg);
619+ #endif
610620 int id = partition.get_local_linear_id ();
611621 T result = sycl::shift_group_left (partition, input, delta);
612622 if ((id + delta) > last_item)
@@ -639,11 +649,21 @@ T shift_sub_group_right(sycl::sub_group sg, T input, int delta, int first_item,
639649 return result;
640650#else
641651 if ((1U << sg.get_local_linear_id ()) & member_mask) {
652+ #if defined(__INTEL_LLVM_COMPILER) && __INTEL_LLVM_COMPILER < 20260000
653+ auto inner = sycl::ext::oneapi::experimental::get_tangle_group (sg);
654+ #else
642655 auto inner = sycl::ext::oneapi::experimental::entangle (sg);
656+ #endif
643657 sycl::group_barrier (inner);
644658 }
659+ #if defined(__INTEL_LLVM_COMPILER) && __INTEL_LLVM_COMPILER < 20260000
660+ auto partition =
661+ sycl::ext::oneapi::experimental::get_fixed_size_group<LogicSubGroupSize>(
662+ sg);
663+ #else
645664 auto partition =
646665 sycl::ext::oneapi::experimental::chunked_partition<LogicSubGroupSize>(sg);
666+ #endif
647667 int id = sg.get_local_linear_id ();
648668 T result = sycl::shift_group_right (partition, input, delta);
649669 if ((id - first_item) < delta)
0 commit comments