[SYCL] Default initialize group_local_memory_for_overwrite (#8472)

In accordance with the extension introducing it, the
group_local_memory_for_overwrite should default initialize its data.
This commit adds this behavior, making it only initialize and
synchronize when needed.

Additionally it changes the local ID calculation in group_local_memory
use the group, as suggested in the TODO.

Signed-off-by: Larsen, Steffen <[email protected]>

Author: steffenlarsen

sycl/include/sycl/ext/oneapi/group_local_memory.hpp

@@ -32,6 +32,12 @@ std::enable_if_t<
 #ifdef __SYCL_DEVICE_ONLY__
   __attribute__((opencl_local)) std::uint8_t *AllocatedMem =
       __sycl_allocateLocalMemory(sizeof(T), alignof(T));
+  // If the type is non-trivial we need to default initialize it.
+  if constexpr (!std::is_trivial_v<T>) {
+    if (g.get_local_linear_id() == 0)
+      new (AllocatedMem) T; // Default initialize.
+    sycl::detail::workGroupBarrier();
+  }
   return reinterpret_cast<__attribute__((opencl_local)) T *>(AllocatedMem);
 #else
   throw feature_not_supported(
@@ -46,19 +52,16 @@ std::enable_if_t<
         sycl::detail::is_group<Group>::value,
     multi_ptr<T, access::address_space::local_space, access::decorated::legacy>>
     __SYCL_ALWAYS_INLINE group_local_memory(Group g, Args &&...args) {
-  (void)g;
 #ifdef __SYCL_DEVICE_ONLY__
   __attribute__((opencl_local)) std::uint8_t *AllocatedMem =
       __sycl_allocateLocalMemory(sizeof(T), alignof(T));
-
-  // TODO switch to using group::get_local_linear_id here once it's implemented
-  id<3> Id = __spirv::initLocalInvocationId<3, id<3>>();
-  if (Id == id<3>(0, 0, 0))
+  if (g.get_local_linear_id() == 0)
     new (AllocatedMem) T(std::forward<Args>(args)...);
   sycl::detail::workGroupBarrier();
   return reinterpret_cast<__attribute__((opencl_local)) T *>(AllocatedMem);
 #else
   // Silence unused variable warning
+  (void)g;
   [&args...] {}();
   throw feature_not_supported(
       "sycl_ext_oneapi_local_memory extension is not supported on host device",