[SYCL] Add a test for generated device code for the group_load_store extension

sycl/include/sycl/ext/oneapi/experimental/group_load_store.hpp

@@ -51,8 +51,15 @@ struct full_group_key
 inline constexpr full_group_key::value_t full_group;
 
 namespace detail {
+struct naive_key : detail::compile_time_property_key<detail::PropKind::Naive> {
+  using value_t = property_value<naive_key>;
+};
+inline constexpr naive_key::value_t naive;
 using namespace sycl::detail;
+} // namespace detail
 
+#ifdef __SYCL_DEVICE_ONLY__
+namespace detail {
 template <typename InputIteratorT, typename OutputElemT>
 inline constexpr bool verify_load_types =
     std::is_same_v<
@@ -101,21 +108,27 @@ int get_mem_idx(GroupTy g, int vec_or_array_idx) {
 }
 } // namespace detail
 
-#ifdef __SYCL_DEVICE_ONLY__
 // Load API span overload.
 template <typename Group, typename InputIteratorT, typename OutputT,
           std::size_t ElementsPerWorkItem,
           typename Properties = decltype(properties())>
 std::enable_if_t<detail::verify_load_types<InputIteratorT, OutputT> &&
                  detail::is_generic_group_v<Group>>
 group_load(Group g, InputIteratorT in_ptr,
-           span<OutputT, ElementsPerWorkItem> out, Properties properties = {}) {
-  constexpr bool blocked = detail::isBlocked(properties);
-
-  group_barrier(g);
-  for (int i = 0; i < out.size(); ++i)
-    out[i] = in_ptr[detail::get_mem_idx<blocked, ElementsPerWorkItem>(g, i)];
-  group_barrier(g);
+           span<OutputT, ElementsPerWorkItem> out, Properties props = {}) {
+  constexpr bool blocked = detail::isBlocked(props);
+
+  if constexpr (props.template has_property<detail::naive_key>()) {
+    group_barrier(g);
+    for (int i = 0; i < out.size(); ++i)
+      out[i] = in_ptr[detail::get_mem_idx<blocked, ElementsPerWorkItem>(g, i)];
+    group_barrier(g);
+  } else {
+    using use_naive =
+        detail::merged_properties_t<Properties,
+                                    decltype(properties(detail::naive))>;
+    return group_load(g, in_ptr, out, use_naive{});
+  }
 }
 
 // Store API span overload.

sycl/include/sycl/ext/oneapi/properties/property.hpp

@@ -199,8 +199,9 @@ enum PropKind : uint32_t {
   DataPlacement = 58,
   ContiguousMemory = 59,
   FullGroup = 60,
+  Naive = 61,
   // PropKindSize must always be the last value.
-  PropKindSize = 61,
+  PropKindSize = 62,
 };
 
 struct property_key_base_tag {};