diff --git a/include/cute/arch/copy_xe.hpp b/include/cute/arch/copy_xe.hpp index d8879bdc3..c05416f9f 100644 --- a/include/cute/arch/copy_xe.hpp +++ b/include/cute/arch/copy_xe.hpp @@ -113,7 +113,7 @@ struct XE_1D_LDSM { copy(const S_ &src, D_ &dst) { #if defined(SYCL_INTEL_TARGET) CUTE_STATIC_ASSERT(sizeof(S_) == sizeof(S)); - auto sg = sycl::ext::oneapi::experimental::this_nd_item<3>().get_sub_group(); + auto sg = sycl::ext::oneapi::this_work_item::get_nd_item<3>().get_sub_group(); *(sycl::vec*)(&dst) = sg.load(sycl::address_space_cast(&*&src)); @@ -183,7 +183,7 @@ struct XE_1D_LOAD_GLOBAL { #if defined(SYCL_INTEL_TARGET) CUTE_STATIC_ASSERT(sizeof(S_) == sizeof(S)); CUTE_STATIC_ASSERT(sizeof(D_) == sizeof(D)); - auto sg = sycl::ext::oneapi::experimental::this_nd_item<3>().get_sub_group(); + auto sg = sycl::ext::oneapi::this_work_item::get_nd_item<3>().get_sub_group(); *(sycl::vec*)(&dst) = sg.load(sycl::address_space_cast(&*&src)); @@ -211,7 +211,7 @@ struct XE_1D_STSM { CUTE_HOST_DEVICE static void copy(S_ const& src, D_ & dst) { #if defined(SYCL_INTEL_TARGET) - auto sg = sycl::ext::oneapi::experimental::this_nd_item<3>().get_sub_group(); + auto sg = sycl::ext::oneapi::this_work_item::get_nd_item<3>().get_sub_group(); sg.store(sycl::address_space_cast(&*&dst), *(sycl::vec*)(&src)); #else @@ -235,7 +235,7 @@ struct XE_1D_STORE_GLOBAL { CUTE_HOST_DEVICE static void copy(S_ const& src, D_ &dst) { #if defined(SYCL_INTEL_TARGET) - auto sg = sycl::ext::oneapi::experimental::this_nd_item<3>().get_sub_group(); + auto sg = sycl::ext::oneapi::this_work_item::get_nd_item<3>().get_sub_group(); sg.store(sycl::address_space_cast(&*&dst), *(sycl::vec*)(&src)); #else diff --git a/include/cutlass/cutlass.h b/include/cutlass/cutlass.h index c9802d840..fbf6276f9 100644 --- a/include/cutlass/cutlass.h +++ b/include/cutlass/cutlass.h @@ -149,14 +149,14 @@ int canonical_warp_group_idx() { #if defined(SYCL_INTEL_TARGET) CUTLASS_DEVICE auto get_sub_group_id() { - return sycl::ext::oneapi::experimental::this_nd_item<3>() + return sycl::ext::oneapi::this_work_item::get_nd_item<3>() .get_sub_group() .get_group_id()[0]; } CUTLASS_DEVICE auto get_sub_group_local_id() { - return sycl::ext::oneapi::experimental::this_nd_item<3>() + return sycl::ext::oneapi::this_work_item::get_nd_item<3>() .get_sub_group() .get_local_id()[0]; }