add test and improve diags

Author: Victor Lomuller

clang/include/clang/Basic/AttrDocs.td

@@ -4096,8 +4096,9 @@ def SYCLWGScopeDocs : Documentation {
   let Heading = "__sycl_detail__::wg_scope";
   let Content = [{
 This attribute can only be applied to records with a trivial default constructor and destructor.
+Types with this attribute cannot be used for non-static data members.
 It indicates that any block and namespace scope variable of a type holding this attribute
-will be allocated in the local memory. For variables allocated in block scope, they behave
+will be allocated in local memory. For variables allocated in block scope, they behave
 as implicitly declared as static.
   }];
 }

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -12466,6 +12466,8 @@ def warn_sycl_kernel_too_big_args : Warning<
 def err_sycl_wg_scope : Error<
   "SYCL work group scope only applies to class with a trivial "
   "%select{default constructor|destructor}0">;
+def err_sycl_field_with_wg_scope : Error<
+  "type with a SYCL work group scope attribute cannot be used with a non-static data members">;
 def err_sycl_virtual_types : Error<
   "no class with a vtable can be used in a SYCL kernel or any code included in the kernel">;
 def note_sycl_recursive_function_declared_here: Note<"function implemented using recursion declared here">;

clang/include/clang/Sema/SemaSYCL.h

@@ -266,7 +266,7 @@ class SemaSYCL : public SemaBase {
 
   void CheckSYCLKernelCall(FunctionDecl *CallerFunc,
                            ArrayRef<const Expr *> Args);
-  void CheckSYCLScope(CXXRecordDecl *Decl);
+  void CheckSYCLScopeAttr(CXXRecordDecl *Decl);
 
   /// Creates a SemaDiagnosticBuilder that emits the diagnostic if the current
   /// context is "used as device code".

clang/lib/Sema/SemaDecl.cpp

@@ -18520,6 +18520,14 @@ FieldDecl *Sema::CheckFieldDecl(DeclarationName Name, QualType T,
     InvalidDecl = true;
   }
 
+  if (LangOpts.SYCLIsDevice) {
+    const CXXRecordDecl *RD = T->getAsCXXRecordDecl();
+    if (RD && RD->hasAttr<SYCLScopeAttr>()) {
+      Diag(Loc, diag::err_sycl_field_with_wg_scope);
+      InvalidDecl = true;
+    }
+  }
+
   if (LangOpts.OpenCL) {
     // OpenCL v1.2 s6.9b,r & OpenCL v2.0 s6.12.5 - The following types cannot be
     // used as structure or union field: image, sampler, event or block types.

clang/lib/Sema/SemaDeclCXX.cpp

@@ -7218,7 +7218,7 @@ void Sema::CheckCompletedCXXClass(Scope *S, CXXRecordDecl *Record) {
       checkCUDADeviceBuiltinTextureClassTemplate(*this, Record);
   }
   if (getLangOpts().SYCLIsDevice && Record->hasAttr<SYCLScopeAttr>()) {
-    SYCL().CheckSYCLScope(Record);
+    SYCL().CheckSYCLScopeAttr(Record);
   }
 }
 

clang/lib/Sema/SemaSYCL.cpp

@@ -5070,11 +5070,14 @@ void SemaSYCL::CheckSYCLKernelCall(FunctionDecl *KernelFunc,
     KernelFunc->setInvalidDecl();
 }
 
-void SemaSYCL::CheckSYCLScope(CXXRecordDecl *Decl) {
+void SemaSYCL::CheckSYCLScopeAttr(CXXRecordDecl *Decl) {
   assert(Decl->hasAttr<SYCLScopeAttr>());
 
   bool HasError = false;
 
+  if (Decl->isDependentContext())
+    return;
+
   // We don't emit both diags at the time as note will only be emitted for the
   // first, which is confusing. So we check both cases but only report one.
   if (!Decl->hasTrivialDefaultConstructor()) {

clang/test/SemaSYCL/sycl_wg_scope.cpp

@@ -0,0 +1,77 @@
+// RUN: %clang_cc1 -fsycl-is-device -verify %s
+
+class [[__sycl_detail__::wg_scope]] G1 {};
+class [[__sycl_detail__::wg_scope]] G2 {
+  G2() = default;
+  G2(int i) : i(i) {}
+  int i;
+};
+
+class [[__sycl_detail__::wg_scope]] G3 {
+  ~G3() = default;
+};
+
+class [[__sycl_detail__::wg_scope]] B4 { // expected-error {{SYCL work group scope only applies to class with a trivial default constructor}}
+  B4() {}
+};
+
+class [[__sycl_detail__::wg_scope]] B5 { // expected-error {{SYCL work group scope only applies to class with a trivial destructor}}
+  ~B5() {}
+};
+
+class [[__sycl_detail__::wg_scope]] B6 { // expected-error {{SYCL work group scope only applies to class with a trivial default constructor}}
+  B6() {}
+  ~B6() {}
+};
+
+template <typename T> class [[__sycl_detail__::wg_scope]] B7 { // #B7
+public:
+  T obj;
+};
+
+struct Valid {};
+struct InvalidCtor {
+  InvalidCtor() {}
+};
+struct InvalidDtor {
+  ~InvalidDtor() {}
+};
+struct InvalidCDtor {
+  InvalidCDtor() {}
+  ~InvalidCDtor() {}
+};
+
+B7<Valid> b7;
+// expected-error@#B7 {{SYCL work group scope only applies to class with a trivial default constructor}}
+// expected-note@+1 {{in instantiation of template class 'B7<InvalidCtor>' requested here}}
+B7<InvalidCtor> b9;
+// expected-error@#B7 {{SYCL work group scope only applies to class with a trivial destructor}}
+// expected-note@+1 {{in instantiation of template class 'B7<InvalidDtor>' requested here}}
+B7<InvalidDtor> b10;
+// expected-error@#B7 {{SYCL work group scope only applies to class with a trivial default constructor}}
+// expected-note@+1 {{in instantiation of template class 'B7<InvalidCDtor>' requested here}}
+B7<InvalidCDtor> b11;
+
+template <typename T> class [[__sycl_detail__::wg_scope]] B12 { // #B12
+public:
+  B12() = default;
+  ~B12() = default;
+  T obj;
+};
+
+B12<Valid> b12;
+// expected-error@#B12 {{SYCL work group scope only applies to class with a trivial default constructor}}
+// expected-note@+1 {{in instantiation of template class 'B12<InvalidCtor>' requested here}}
+B12<InvalidCtor> b13;
+
+class B14 {
+  G1 field; // expected-error {{type with a SYCL work group scope attribute cannot be used with a non-static data members}}
+};
+
+template <typename T> class B15 {
+  T field; // #B15-field
+};
+
+// expected-error@#B15-field {{type with a SYCL work group scope attribute cannot be used with a non-static data members}}
+// expected-note@+1 {{in instantiation of template class 'B15<G1>' requested here}}
+B15<G1> b15;

llvm/lib/SYCLLowerIR/LowerWGLocalMemory.cpp

@@ -129,20 +129,17 @@ lowerDynamicLocalMemCallDirect(CallInst *CI, Triple TT,
 
   Value *GVPtr = [&]() -> Value * {
     IRBuilder<> Builder(CI);
-    if (TT.isSPIROrSPIRV()) {
-
+    if (TT.isSPIROrSPIRV())
       return Builder.CreateLoad(CI->getType(), LocalMemPlaceholder);
-    } else {
-      Value *ArgAlign = CI->getArgOperand(0);
-      Align RequestedAlignment{
-          cast<llvm::ConstantInt>(ArgAlign)->getZExtValue()};
-      MaybeAlign CurrentAlignment = LocalMemPlaceholder->getAlign();
-      if (!CurrentAlignment.has_value() ||
-          (CurrentAlignment.value() < RequestedAlignment))
-        LocalMemPlaceholder->setAlignment(RequestedAlignment);
-
-      return Builder.CreatePointerCast(LocalMemPlaceholder, CI->getType());
-    }
+    Value *ArgAlign = CI->getArgOperand(0);
+    Align RequestedAlignment{
+        cast<llvm::ConstantInt>(ArgAlign)->getZExtValue()};
+    MaybeAlign CurrentAlignment = LocalMemPlaceholder->getAlign();
+    if (!CurrentAlignment.has_value() ||
+        (CurrentAlignment.value() < RequestedAlignment))
+      LocalMemPlaceholder->setAlignment(RequestedAlignment);
+
+    return Builder.CreatePointerCast(LocalMemPlaceholder, CI->getType());
   }();
   CI->replaceAllUsesWith(GVPtr);
 }

llvm/test/SYCLLowerIR/work_group_static.ll

@@ -0,0 +1,42 @@
+; RUN: opt -S -sycllowerwglocalmemory -bugpoint-enable-legacy-pm < %s | FileCheck %s
+; RUN: opt -S -passes=sycllowerwglocalmemory < %s | FileCheck %s
+
+; CHECK-NOT: __sycl_dynamicLocalMemoryPlaceholder
+
+target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64"
+target triple = "spir64-unknown-unknown"
+
+; CHECK-DAG: @__sycl_dynamicLocalMemoryPlaceholder_GV = linkonce_odr addrspace(3) global ptr addrspace(3) undef
+
+; Function Attrs: convergent norecurse
+; CHECK-DAG: @_ZTS7KernelA(ptr addrspace(1) %0, ptr addrspace(3) noalias "sycl-implicit-local-arg" %[[IMPLICT_ARG:[a-zA-Z0-9]+]]{{.*}} !kernel_arg_addr_space ![[ADDR_SPACE_MD:[0-9]+]]
+define weak_odr dso_local spir_kernel void @_ZTS7KernelA(ptr addrspace(1) %0) local_unnamed_addr #0 !kernel_arg_addr_space !5 {
+entry:
+; CHECK-DAG: store ptr addrspace(3) %[[IMPLICT_ARG]], ptr addrspace(3) @__sycl_dynamicLocalMemoryPlaceholder_GV
+; CHECK-DAG: %[[LD1:[a-zA-Z0-9]+]] = load ptr addrspace(3), ptr addrspace(3) @__sycl_dynamicLocalMemoryPlaceholder_GV
+  %1 = tail call spir_func ptr addrspace(3) @__sycl_dynamicLocalMemoryPlaceholder(i64 128) #1
+; CHECK-DAG: getelementptr inbounds i8, ptr addrspace(3) %[[LD1]]
+  %2 = getelementptr inbounds i8, ptr addrspace(3) %1, i64 4
+; CHECK-DAG: %[[LD2:[a-zA-Z0-9]+]] = load ptr addrspace(3), ptr addrspace(3) @__sycl_dynamicLocalMemoryPlaceholder_GV
+  %3 = tail call spir_func ptr addrspace(3) @__sycl_dynamicLocalMemoryPlaceholder(i64 4) #1
+  ret void
+}
+
+; Function Attrs: convergent
+declare dso_local spir_func ptr addrspace(3) @__sycl_dynamicLocalMemoryPlaceholder(i64) local_unnamed_addr #1
+
+attributes #0 = { convergent norecurse "disable-tail-calls"="false" "frame-pointer"="all" "less-precise-fpmad"="false" "min-legal-vector-width"="0" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size"="true" "unsafe-fp-math"="false" "use-soft-float"="false" "sycl-work-group-static"="1" }
+attributes #1 = { convergent norecurse }
+
+!llvm.module.flags = !{!0}
+!opencl.spir.version = !{!1}
+!spirv.Source = !{!2}
+!llvm.ident = !{!3}
+
+!0 = !{i32 1, !"wchar_size", i32 4}
+!1 = !{i32 1, i32 2}
+!2 = !{i32 4, i32 100000}
+!3 = !{!"clang version 13.0.0"}
+!4 = !{}
+; ![[ADDR_SPACE_MD]] = !{i32 1, i32 3}
+!5 = !{i32 1}

sycl/doc/extensions/proposed/sycl_ext_oneapi_work_group_static.asciidoc

@@ -64,7 +64,7 @@ Device local memory is memory that is shared by all work-items in a work-group.
 The behavior is similar to the CUDA `+__shared__+` keyword, and the extension
 draws some inspiration from the {cpp} `thread_local` keyword.
 
-`work_group_static` variables can be allocated at namespace or block scope,
+`work_group_static` can only be used to declare variables at namespace, block or class scope,
 lifting many of the restrictions in the existing
 link:../supported/sycl_ext_oneapi_local_memory.asciidoc[sycl_ext_oneapi_local_memory]
 extension. Note, however, that `work_group_static` variables currently place
@@ -106,7 +106,7 @@ an object into device local memory.
 namespace sycl::ext::oneapi::experimental {
 
 template <typename T>
-class work_group_static {
+class work_group_static final {
 public:
 
   work_group_static() = default;
@@ -121,8 +121,7 @@ public:
   T* operator&() const noexcept;
 
 private:
-  T* ptr; // exposition only
-
+  T storage;
 };
 
 } // namespace sycl::ext::oneapi::experimental
@@ -134,6 +133,9 @@ The storage for the object is allocated in device local memory before
 calling the user's kernel lambda, and deallocated when all work-items
 in the work-group have completed execution of the kernel.
 
+Objects of type `work_group_static` must only be declared at namespace, block, lambda or class scope.
+If the object is declared in class scope, it must be declared as a static data member.
+
 SYCL implementations conforming to the full feature set treat
 `work_group_static` similarly to the `thread_local` keyword, and when
 a `work_group_static` object is declared at block scope it behaves
@@ -150,18 +152,11 @@ multiple times, developers must take care to avoid race conditions (e.g., by
 calling `group_barrier` before and after using the memory).
 ====
 
-Change to SYCL 2020 section `5.9.2 Common address space deduction rules`:
-Namespace scope: if the variable is `work_group_static` object,
-then the variable is assigned to the local address space.
-Otherwise normal rules applies.
-
 SYCL 2020 requires that all global variables accessed by a device function are
 `const` or `constexpr`. This extension lifts that restriction for
 `work_group_static` variables.
 
-When `T` is a class type or bounded array, the size of the allocation is known
-at compile-time, and a SYCL implementation embeds the size of the allocation
-directly within a kernel. Each instance of `work_group_static<T>` is associated
+Each instance of `work_group_static<T>` is associated
 with a unique allocation in device local memory.
 
 [source,c++]
@@ -173,7 +168,7 @@ associated with this instance of `work_group_static`.
 
 [source,c++]
 ----
-work_group_static<T>& operator=(const T& value) noexcept;
+work_group_static& operator=(const T& value) noexcept;
 ----
 _Constraints_: Available only if `std::is_array_v<T>` is false.
 
@@ -188,6 +183,11 @@ T* operator&() noexcept;
 _Returns_: A pointer to the device local memory associated with this
 instance of `work_group_static` (i.e., `ptr`).
 
+==== Interaction with common address space deduction rules
+
+Objects of type `work_group_static` are assigned to
+the local address space.
+
 === `get_dynamic_work_group_memory` function
 
 The `get_dynamic_work_group_memory` function provides access
@@ -213,15 +213,7 @@ in device local memory, regardless of `T`. For example, two call declared
 as `get_dynamic_work_group_memory<int>` and
 `get_dynamic_work_group_memory<float>` will be associated with the same shared allocation.
 
-If the total amount of device local memory requested (i.e., the sum of
-all memory requested by `local_accessor`, `group_local_memory`,
-`group_local_memory_for_overwrite` and `work_group_static`) exceeds a device's
-local memory capacity (as reported by `local_mem_size`) then the implementation
-must throw a synchronous `exception` with the `errc::memory_allocation` error
-code from the kernel invocation command (e.g. `parallel_for`).
-
-
-==== Kernel properties
+=== Kernel properties
 
 The `work_group_static_size` property must be passed to a kernel to determine
 the run-time size of the device local memory allocation associated with
@@ -252,6 +244,14 @@ device local memory required by the kernel in bytes.
 
 |===
 
+=== Total allocation check
+
+If the total amount of device local memory requested (i.e., the sum of
+all memory requested by `local_accessor`, `group_local_memory`,
+`group_local_memory_for_overwrite`, `work_group_static` and `work_group_static_size`) exceeds a device's
+local memory capacity (as reported by `local_mem_size`) then the implementation
+must throw a synchronous `exception` with the `errc::memory_allocation` error
+code from the kernel invocation command (e.g. `parallel_for`).
 
 ==== Usage examples