Add support for unbounded arrays

Author: lbushi25

clang/lib/CodeGen/CodeGenTypes.cpp

@@ -488,6 +488,14 @@ llvm::Type *CodeGenTypes::ConvertType(QualType T) {
         return Ty;
     }
   }
+  // An incomplete array AST type is typically lowered to an array of length zero in LLVM IR.
+  // For SYCL devices, this is incompatible with SPIRV which does not accept arrays of length zero
+  // so we explicitly intercept this case to instead lower to an array of length 1 instead.
+  if (Context.getLangOpts().SYCLIsDevice) 
+	if (T->isIncompleteArrayType()) {
+		return llvm::ArrayType::get(ConvertType(cast<ArrayType>(T)->getElementType()), 1);
+}
+}
 
   // RecordTypes are cached and processed specially.
   if (const RecordType *RT = dyn_cast<RecordType>(Ty))

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

@@ -22,19 +22,17 @@ inline constexpr bool is_unbounded_array_v = is_unbounded_array<T>::value;
 
 class work_group_memory_impl {
 public:
-  work_group_memory_impl() : wgm_size{0}, buffer_size{0} {}
+  work_group_memory_impl() : buffer_size{0} {}
   work_group_memory_impl(const work_group_memory_impl &rhs) = default;
   work_group_memory_impl &
   operator=(const work_group_memory_impl &rhs) = default;
-  work_group_memory_impl(size_t wgm_size, size_t buffer_size)
-      : wgm_size{wgm_size}, buffer_size{buffer_size} {}
-  size_t wgm_size;
+  work_group_memory_impl(size_t buffer_size)
+      : buffer_size{buffer_size} {}
+  private:
   size_t buffer_size;
+  friend class sycl::handler;
 };
 
-inline size_t getWorkGroupMemoryOwnSize(detail::work_group_memory_impl *wgm) {
-  return wgm->wgm_size;
-}
 } // namespace detail
 
 namespace ext::oneapi::experimental {
@@ -55,13 +53,12 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(work_group_memory) work_group_memory
   template <typename T = DataT,
             typename = std::enable_if_t<!sycl::detail::is_unbounded_array_v<T>>>
   work_group_memory(handler &)
-      : sycl::detail::work_group_memory_impl(sizeof(work_group_memory),
+      : sycl::detail::work_group_memory_impl(
                                              sizeof(DataT)) {}
   template <typename T = DataT,
             typename = std::enable_if_t<sycl::detail::is_unbounded_array_v<T>>>
   work_group_memory(size_t num, handler &)
       : sycl::detail::work_group_memory_impl(
-            sizeof(work_group_memory),
             num * sizeof(std::remove_extent_t<DataT>)) {}
   template <access::decorated IsDecorated = access::decorated::no>
   multi_ptr<value_type, access::address_space::local_space, IsDecorated>
@@ -70,7 +67,7 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(work_group_memory) work_group_memory
                                     IsDecorated, value_type>(ptr);
   }
   DataT *operator&() const { return reinterpret_cast<DataT *>(ptr); }
-  operator DataT &() const { return *reinterpret_cast<DataT *>(ptr); }
+  operator DataT &() const { return *(this->operator&()); }
   template <typename T = DataT,
             typename = std::enable_if_t<!std::is_array_v<T>>>
   const work_group_memory &operator=(const DataT &value) const {

sycl/include/sycl/handler.hpp

@@ -171,8 +171,6 @@ class graph_impl;
 } // namespace ext::oneapi::experimental::detail
 namespace detail {
 
-class work_group_memory_impl;
-size_t getWorkGroupMemoryOwnSize(work_group_memory_impl *);
 class handler_impl;
 class kernel_impl;
 class queue_impl;
@@ -720,8 +718,7 @@ class __SYCL_EXPORT handler {
       const ext::oneapi::experimental::work_group_memory<DataT, PropertyListT>
           &Arg) {
     addArg(detail::kernel_param_kind_t::kind_work_group_memory, &Arg,
-           detail::getWorkGroupMemoryOwnSize(
-               static_cast<detail::work_group_memory_impl *>(&Arg)),
+           0,
            ArgIndex);
   }
 

sycl/test-e2e/WorkGroupMemory/swap_test.cpp

@@ -186,6 +186,29 @@ void swap_array_1d(T (&a)[N], T (&b)[N], size_t batch_size) {
   for (int i = 0; i < N; ++i) {
     assert(a[i] == old_b[i] && b[i] == old_a[i] && "Incorrect swap!");
   }
+
+  // Same as above but use an unbounded array as temporary storage
+  {
+    sycl::buffer<T, 1> buf_a{a, N};
+    sycl::buffer<T, 1> buf_b{b, N};
+    q.submit([&](sycl::handler &cgh) {
+      sycl::accessor acc_a{buf_a, cgh};
+      sycl::accessor acc_b{buf_b, cgh};
+      syclexp::work_group_memory<T[]> temp{N, cgh};
+      sycl::nd_range<1> ndr{size, wgsize};
+      cgh.parallel_for(ndr, [=](sycl::nd_item<> it) {
+        const auto i = it.get_global_id();
+        auto ptr = &temp;
+        (*ptr)[i] = acc_a[i];
+        acc_a[i] = acc_b[i];
+        acc_b[i] = (*ptr)[i];
+      });
+    });
+  }
+  // Four swaps same as no swap
+  for (int i = 0; i < N; ++i) {
+    assert(a[i] == old_a[i] && b[i] == old_b[i] && "Incorrect swap!");
+  }
 }
 
 template <typename T, size_t N>
@@ -283,6 +306,33 @@ void swap_array_2d(T (&a)[N][N], T (&b)[N][N], size_t batch_size) {
              "Incorrect swap!");
     }
   }
+  
+  // Same as above but use an unbounded array as temporary storage
+{
+    sycl::buffer<T, 2> buf_a{a[0], sycl::range{N, N}};
+    sycl::buffer<T, 2> buf_b{b[0], sycl::range{N, N}};
+    q.submit([&](sycl::handler &cgh) {
+      sycl::accessor acc_a{buf_a, cgh};
+      sycl::accessor acc_b{buf_b, cgh};
+      syclexp::work_group_memory<T[][N]> temp{N, cgh};
+      sycl::nd_range<2> ndr{size, wgsize};
+      cgh.parallel_for(ndr, [=](sycl::nd_item<2> it) {
+        const auto i = it.get_global_id()[0];
+        const auto j = it.get_global_id()[1];
+        temp[i][j] = acc_a[i][j];
+        acc_a[i][j] = acc_b[i][j];
+        syclexp::work_group_memory<T[][N]> temp2{temp};
+        acc_b[i][j] = temp2[i][j];
+      });
+    });
+  } 
+  for (int i = 0; i < N; ++i) {
+    for (int j = 0; j < N; ++j) {
+      // Four swaps are the same as no swap
+      assert(a[i][j] == old_a[i][j] && b[i][j] == old_b[i][j] &&
+             "Incorrect swap!");
+    }
+  }
 }
 
 // Coherency test that checks that work group memory is truly shared by