More robust checking for fence variable signaling extension is present

Author: oleksandr-pavlyk

dpctl/tensor/libtensor/include/utils/sycl_utils.hpp

@@ -423,7 +423,8 @@ struct Identity<Op, T, std::enable_if_t<UseBuiltInIdentity<Op, T>::value>>
 
 // Sub-group load/store
 
-#if defined(SYCL_EXT_ONEAPI_GROUP_LOAD_STORE)
+#if defined(SYCL_EXT_ONEAPI_GROUP_LOAD_STORE) &&                               \
+    (SYCL_EXT_ONEAPI_GROUP_LOAD_STORE)
 namespace ls_ns = sycl::ext::oneapi::experimental;
 #endif
 
@@ -434,7 +435,8 @@ template <std::uint8_t vec_sz,
 auto sub_group_load(const sycl::sub_group &sg,
                     sycl::multi_ptr<ElementType, Space, DecorateAddress> m_ptr)
 {
-#if defined(SYCL_EXT_ONEAPI_GROUP_LOAD_STORE)
+#if defined(SYCL_EXT_ONEAPI_GROUP_LOAD_STORE) &&                               \
+    (SYCL_EXT_ONEAPI_GROUP_LOAD_STORE)
     sycl::vec<ElementType, vec_sz> x;
     ls_ns::group_load(sg, m_ptr, x);
     return x;
@@ -449,7 +451,8 @@ template <sycl::access::address_space Space,
 auto sub_group_load(const sycl::sub_group &sg,
                     sycl::multi_ptr<ElementType, Space, DecorateAddress> m_ptr)
 {
-#if defined(SYCL_EXT_ONEAPI_GROUP_LOAD_STORE)
+#if defined(SYCL_EXT_ONEAPI_GROUP_LOAD_STORE) &&                               \
+    (SYCL_EXT_ONEAPI_GROUP_LOAD_STORE)
     ElementType x;
     ls_ns::group_load(sg, m_ptr, x);
     return x;
@@ -466,10 +469,13 @@ void sub_group_store(const sycl::sub_group &sg,
                      const sycl::vec<ElementType, vec_sz> &val,
                      sycl::multi_ptr<ElementType, Space, DecorateAddress> m_ptr)
 {
-#if defined(SYCL_EXT_ONEAPI_GROUP_LOAD_STORE)
+#if defined(SYCL_EXT_ONEAPI_GROUP_LOAD_STORE) &&                               \
+    (SYCL_EXT_ONEAPI_GROUP_LOAD_STORE)
     ls_ns::group_store(sg, val, m_ptr);
+    return;
 #else
     sg.store<vec_sz>(m_ptr, val);
+    return;
 #endif
 }
 
@@ -480,10 +486,13 @@ void sub_group_store(const sycl::sub_group &sg,
                      const ElementType &val,
                      sycl::multi_ptr<ElementType, Space, DecorateAddress> m_ptr)
 {
-#if defined(SYCL_EXT_ONEAPI_GROUP_LOAD_STORE)
+#if defined(SYCL_EXT_ONEAPI_GROUP_LOAD_STORE) &&                               \
+    (SYCL_EXT_ONEAPI_GROUP_LOAD_STORE)
     ls_ns::group_store(sg, val, m_ptr);
+    return;
 #else
     sg.store(m_ptr, val);
+    return;
 #endif
 }