Skip to content

Commit 0e99dc7

Browse files
authored
[SYCLomatic] Backwards compatibility with older oneAPI compilers for non-uniform group APIs (#2997)
Signed-off-by: Matthew Michel <matthew.michel@intel.com>
1 parent f350c9d commit 0e99dc7

1 file changed

Lines changed: 20 additions & 0 deletions

File tree

  • clang/runtime/dpct-rt/include/dpct

clang/runtime/dpct-rt/include/dpct/util.hpp

Lines changed: 20 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -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

Comments
 (0)