Replace deprecated this_nd_item calls (#150)

codeplaysoftware · Oct 28, 2024 · 641f717 · 641f717
1 parent 1dc136d
commit 641f717
Show file tree

Hide file tree

Showing 2 changed files with 6 additions and 6 deletions.
diff --git 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<S_, N>*)(&dst)
         = sg.load<N>(sycl::address_space_cast<sycl::access::address_space::local_space,
                   sycl::access::decorated::yes>(&*&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<S_, N>*)(&dst) 
         = sg.load<N>(sycl::address_space_cast<sycl::access::address_space::global_space,
                   sycl::access::decorated::yes>(&*&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<N>(sycl::address_space_cast<sycl::access::address_space::local_space,
             sycl::access::decorated::yes>(&*&dst), *(sycl::vec<D_, N>*)(&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<N>(sycl::address_space_cast<sycl::access::address_space::global_space,
             sycl::access::decorated::yes>(&*&dst), *(sycl::vec<D_, N>*)(&src));
     #else

diff --git 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];
 }