sycl backend: Small changes.

fweig · May 22, 2024 · 9068753 · 9068753
1 parent f444797
commit 9068753
Show file tree

Hide file tree

Showing 2 changed files with 7 additions and 3 deletions.
diff --git a/src/xpu/detail/platform/sycl/device.h b/src/xpu/detail/platform/sycl/device.h
@@ -380,13 +380,17 @@ struct xpu::detail::action_runner<xpu::detail::kernel_tag, K, void(K::*)(xpu::ke
             auto cmem_accessors = cmem_traits.make_accessors(cmem_buffers, cgh);
             constants cmem{internal_ctor, cmem_accessors};
 
+            #if defined(__INTEL_LLVM_COMPILER) && __INTEL_LLVM_COMPILER < 20240000
             sycl::stream out{0, 0, cgh};
+            #endif
             cgh.parallel_for<K>(sycl::nd_range<3>{global_range, local_range}, [=](sycl::nd_item<3> item) {
-                // WTF: icpx sometimes optimizes out the kernel call (when using O2)
+                #if defined(__INTEL_LLVM_COMPILER) && __INTEL_LLVM_COMPILER < 20240000
+                // WTF: old versions of icpx sometimes optimizes out the kernel call (when using O2)
                 // if we dont add the print statement
                 if (item.get_global_id(0) == static_cast<size_t>(-1)) {
                     out << "";
                 }
+                #endif
                 shared_memory &smem = shared_memory_acc;
                 tpos pos{internal_ctor, item};
                 context ctx{internal_ctor, pos, smem, cmem};

diff --git a/src/xpu/detail/platform/sycl/tpos_impl.h b/src/xpu/detail/platform/sycl/tpos_impl.h
@@ -8,7 +8,7 @@ namespace xpu::detail {
 class tpos_impl {
 
 public:
-    tpos_impl(sycl::nd_item<3> nd_item) : m_nd_item(nd_item) {}
+    explicit tpos_impl(sycl::nd_item<3> nd_item) : m_nd_item(nd_item) {}
 
     int thread_idx_x() const { return m_nd_item.get_local_id(0); }
     int thread_idx_y() const { return m_nd_item.get_local_id(1); }
@@ -26,7 +26,7 @@ class tpos_impl {
     int grid_dim_y() const { return m_nd_item.get_group_range(1); }
     int grid_dim_z() const { return m_nd_item.get_group_range(2); }
 
-    void barrier() const { m_nd_item.barrier(sycl::access::fence_space::local_space); }
+    void barrier() const { sycl::group_barrier(m_nd_item.get_group()); }
 
     sycl::group<3> group() const { return m_nd_item.get_group(); }