Merge branch 'sycl' into throw_get_profiling_info_exception

Author: HabKaffee

clang/include/clang/Basic/Attr.td

@@ -3039,9 +3039,12 @@ def WorkGroupSizeHint :  InheritableAttr {
   let Spellings = [GNU<"work_group_size_hint">,
                    CXX11<"sycl", "work_group_size_hint">];
   let Args = [ExprArgument<"XDim">,
-              ExprArgument<"YDim">,
-              ExprArgument<"ZDim">];
+              ExprArgument<"YDim", /*optional*/1>,
+              ExprArgument<"ZDim", /*optional*/1>];
   let Subjects = SubjectList<[Function], ErrorDiag>;
+  // In OpenCL C this attribute takes no default values whereas in SYCL it does. 
+  // To avoid confusing diagnostics, the checks are deferred to "handleWorkGroupSizeHint". 
+  let HasCustomParsing = 1;
   let AdditionalMembers = [{
     Optional<llvm::APSInt> getXDimVal() const {
       if (const auto *CE = dyn_cast<ConstantExpr>(getXDim()))

clang/include/clang/Basic/AttrDocs.td

@@ -2848,77 +2848,23 @@ The GNU spelling is deprecated in SYCL mode.
     [[sycl::work_group_size_hint(2, 2, 2)]] void operator()() const {}
   };
 
-The arguments to ``reqd_work_group_size`` are ordered based on which index
+The arguments to ``work_group_size_hint`` are ordered based on which index
 increments the fastest. In OpenCL, the first argument is the index that
 increments the fastest, and in SYCL, the last argument is the index that
 increments the fastest.
 
-In OpenCL, all three arguments are required.
+In OpenCL C, this attribute is available with the GNU spelling
+(``__attribute__((work_group_size_hint(X, Y, Z)))``) and all 
+three arguments are required.
 
 In SYCL, the attribute accepts either one, two, or three arguments; in each
 form, the last (or only) argument is the index that increments fastest. The
 number of arguments passed to the attribute must match the dimensionality of
 the kernel the attribute is applied to.
 
-If the ``reqd_work_group_size attribute`` is specified on a declaration along
-with ``num_simd_work_items``, the required work group size specified by
-``num_simd_work_items`` must evenly divide the index that increments fastest
-in the ``reqd_work_group_size`` attribute.
-
-.. code-block:: c++
-
-  // Note, '64' is evenly divisible by '4'; in SYCL, the last
-  // argument to the attribute is the one which increments fastest.
-  struct func {
-    [[intel::num_simd_work_items(4)]]
-    [[sycl::reqd_work_group_size(7, 4, 64)]]
-    void operator()() const {}
-  };
-
-  // Note, '8' is evenly divisible by '8'; in SYCL, the last
-  // argument to the attribute is the one which increments fastest.
-  struct bar {
-    [[sycl::reqd_work_group_size(1, 1, 8)]]
-    [[intel::num_simd_work_items(8)]]
-    void operator()() const {}
-  };
-
-  // Note, '10' is evenly divisible by '5'; in SYCL, the last
-  // argument to the attribute is the one which increments fastest.
-  [[cl::reqd_work_group_size(7, 5, 10)]]
-  [[intel::num_simd_work_items(5)]] void fun2() {}
-
-  // Note, '8' is evenly divisible by '4'; in SYCL, the last
-  // argument to the attribute is the one which increments fastest.
-  [[intel::num_simd_work_items(4)]]
-  [[cl::reqd_work_group_size(5, 4, 8)]] void fun3() {}
-
-  // Note, '8' is evenly divisible by '8'; in SYCL, the last
-  // argument to the attribute is the one which increments fastest.
-  struct func1 {
-    [[intel::num_simd_work_items(8)]]
-    [[cl::reqd_work_group_size(1, 1, 8)]]
-    void operator()() const {}
-  };
-  
-  // Note, '8' is evenly divisible by '4'; in SYCL, the last
-  // argument to the attribute is the one which increments fastest.
-  struct bar1 {
-    [[cl::reqd_work_group_size(7, 4, 8)]]
-    [[intel::num_simd_work_items(4)]]
-    void operator()() const {}
-  };
-
-  // Note, '4' is evenly divisible by '2'; in SYCL, the last
-  // argument to the attribute is the one which increments fastest.
-  [[intel::num_simd_work_items(2)]]
-  __attribute__((reqd_work_group_size(3, 2, 4))) void test();
- 
-  // Note, '8' is evenly divisible by '2'; in SYCL, the last
-  // argument to the attribute is the one which increments fastest.
-  __attribute__((reqd_work_group_size(3, 2, 8)))
-  [intel::num_simd_work_items(2)]] void test();
-
+In SYCL 1.2.1 mode, the ``sycl::work_group_size_hint`` attribute is propagated 
+from the function it is applied to onto the kernel which calls the function. 
+In SYCL 2020 mode, the attribute is not propagated to the kernel.
   }];
 }
 

clang/lib/Sema/SemaDeclAttr.cpp

@@ -3510,8 +3510,32 @@ Sema::MergeWorkGroupSizeHintAttr(Decl *D, const WorkGroupSizeHintAttr &A) {
 static void handleWorkGroupSizeHint(Sema &S, Decl *D, const ParsedAttr &AL) {
   S.CheckDeprecatedSYCLAttributeSpelling(AL);
 
-  S.AddWorkGroupSizeHintAttr(D, AL, AL.getArgAsExpr(0), AL.getArgAsExpr(1),
-                             AL.getArgAsExpr(2));
+  // __attribute__((work_group_size_hint) requires exactly three arguments.
+  if (AL.getSyntax() == ParsedAttr::AS_GNU || !AL.hasScope() ||
+      (AL.hasScope() && !AL.getScopeName()->isStr("sycl"))) {
+    if (!AL.checkExactlyNumArgs(S, 3))
+      return;
+  }
+
+  // FIXME: NumArgs checking is disabled in Attr.td to keep consistent
+  // disgnostics with OpenCL C that does not have optional values here.
+  if (!AL.checkAtLeastNumArgs(S, 1) || !AL.checkAtMostNumArgs(S, 3))
+    return;
+
+  // Handles default arguments in [[sycl::work_group_size_hint]] attribute.
+  auto SetDefaultValue = [](Sema &S, const ParsedAttr &AL) {
+    assert(AL.getKind() == ParsedAttr::AT_WorkGroupSizeHint && AL.hasScope() &&
+           AL.getScopeName()->isStr("sycl"));
+    return IntegerLiteral::Create(S.Context, llvm::APInt(32, 1),
+                                  S.Context.IntTy, AL.getLoc());
+  };
+
+  Expr *XDimExpr = AL.getArgAsExpr(0);
+  Expr *YDimExpr =
+      AL.isArgExpr(1) ? AL.getArgAsExpr(1) : SetDefaultValue(S, AL);
+  Expr *ZDimExpr =
+      AL.isArgExpr(2) ? AL.getArgAsExpr(2) : SetDefaultValue(S, AL);
+  S.AddWorkGroupSizeHintAttr(D, AL, XDimExpr, YDimExpr, ZDimExpr);
 }
 
 // Handles max_work_group_size attribute.

clang/lib/Sema/SemaSYCL.cpp

@@ -561,8 +561,8 @@ static void collectSYCLAttributes(Sema &S, FunctionDecl *FD,
     llvm::copy_if(FD->getAttrs(), std::back_inserter(Attrs), [](Attr *A) {
       // FIXME: Make this list self-adapt as new SYCL attributes are added.
       return isa<IntelReqdSubGroupSizeAttr, IntelNamedSubGroupSizeAttr,
-                 ReqdWorkGroupSizeAttr, SYCLIntelKernelArgsRestrictAttr,
-                 SYCLIntelNumSimdWorkItemsAttr,
+                 ReqdWorkGroupSizeAttr, WorkGroupSizeHintAttr,
+                 SYCLIntelKernelArgsRestrictAttr, SYCLIntelNumSimdWorkItemsAttr,
                  SYCLIntelSchedulerTargetFmaxMhzAttr,
                  SYCLIntelMaxWorkGroupSizeAttr, SYCLIntelMaxGlobalWorkDimAttr,
                  SYCLIntelNoGlobalWorkOffsetAttr, SYCLSimdAttr>(A);
@@ -3927,6 +3927,22 @@ static void PropagateAndDiagnoseDeviceAttr(
     }
     break;
   }
+  case attr::Kind::WorkGroupSizeHint: {
+    auto *WGSH = cast<WorkGroupSizeHintAttr>(A);
+    if (auto *Existing = SYCLKernel->getAttr<WorkGroupSizeHintAttr>()) {
+      if (Existing->getXDimVal() != WGSH->getXDimVal() ||
+          Existing->getYDimVal() != WGSH->getYDimVal() ||
+          Existing->getZDimVal() != WGSH->getZDimVal()) {
+        S.Diag(SYCLKernel->getLocation(),
+               diag::err_conflicting_sycl_kernel_attributes);
+        S.Diag(Existing->getLocation(), diag::note_conflicting_attribute);
+        S.Diag(WGSH->getLocation(), diag::note_conflicting_attribute);
+        SYCLKernel->setInvalidDecl();
+      }
+    }
+    SYCLKernel->addAttr(A);
+    break;
+  }
   case attr::Kind::SYCLIntelMaxWorkGroupSize: {
     auto *SIMWGSA = cast<SYCLIntelMaxWorkGroupSizeAttr>(A);
     if (auto *Existing = SYCLKernel->getAttr<ReqdWorkGroupSizeAttr>()) {

clang/test/SemaOpenCL/invalid-kernel-attrs.cl

@@ -12,6 +12,8 @@ kernel __attribute__((vec_type_hint(int))) __attribute__((vec_type_hint(float)))
 
 kernel __attribute__((work_group_size_hint(8,16,32,4))) void kernel6() {} //expected-error{{'work_group_size_hint' attribute requires exactly 3 arguments}}
 
+kernel __attribute__((work_group_size_hint(1,2))) void kernel6b() {}  //expected-error{{'work_group_size_hint' attribute requires exactly 3 arguments}}
+
 kernel __attribute__((work_group_size_hint(1,2,3))) __attribute__((work_group_size_hint(3,2,1))) void kernel7() {}  //expected-warning{{attribute 'work_group_size_hint' is already applied with different arguments}} expected-note {{previous attribute is here}}
 
 __attribute__((reqd_work_group_size(8,16,32))) void kernel8(){} // expected-error {{attribute 'reqd_work_group_size' can only be applied to an OpenCL kernel}}

clang/test/SemaSYCL/check-work-group-size-hint-device.cpp

@@ -0,0 +1,172 @@
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -fsyntax-only -sycl-std=2017 -Wno-sycl-2017-compat -verify -DEXPECT_PROP -DTRIGGER_ERROR %s
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2017 -Wno-sycl-2017-compat -ast-dump -DEXPECT_PROP %s | FileCheck %s
+
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -fsyntax-only -sycl-std=2020 -verify -DTRIGGER_ERROR %s
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2020 -ast-dump %s | FileCheck %s
+
+// Test for AST of work_group_size_hint kernel attribute in SYCL 1.2.1. and SYCL 2020 modes. 
+#include "sycl.hpp"
+
+// Check the basics.
+#ifdef TRIGGER_ERROR
+[[sycl::work_group_size_hint]] void f0();                 // expected-error {{'work_group_size_hint' attribute takes at least 1 argument}}
+[[sycl::work_group_size_hint(12, 12, 12, 12)]] void f1(); // expected-error {{'work_group_size_hint' attribute takes no more than 3 arguments}}
+[[sycl::work_group_size_hint("derp", 1, 2)]] void f2();   // expected-error {{integral constant expression must have integral or unscoped enumeration type, not 'const char[5]'}}
+[[sycl::work_group_size_hint(1, 1, 1)]] int i;            // expected-error {{'work_group_size_hint' attribute only applies to functions}}
+#endif
+
+// Produce a conflicting attribute warning when the args are different.
+[[sycl::work_group_size_hint(4, 1, 1)]] void f3();    // expected-note {{previous attribute is here}}
+[[sycl::work_group_size_hint(1, 1, 32)]] void f3() {} // expected-warning {{attribute 'work_group_size_hint' is already applied with different arguments}}
+
+// 1 and 2 dim versions
+[[sycl::work_group_size_hint(2)]] void f4();    // ok
+[[sycl::work_group_size_hint(2, 1)]] void f5(); // ok
+
+// FIXME: This turns out to be wrong as there aren't really default values
+// (that is an implementation detail we use but shouldn't expose to the user).
+// Instead, the dimensionality of the attribute needs to match that of the
+// kernel, so the one, two, and three arg forms of the attribute are actually
+// *different* attributes. This means that you should not be able to redeclare
+// the function with a different dimensionality.
+// As a result these two (re)declarations should result in errors.
+[[sycl::work_group_size_hint(2)]] void f5();
+[[sycl::work_group_size_hint(2, 1, 1)]] void f5();
+
+// Catch the easy case where the attributes are all specified at once with
+// different arguments.
+[[sycl::work_group_size_hint(4, 1, 1), sycl::work_group_size_hint(32, 1, 1)]] void f7(); // expected-warning {{attribute 'work_group_size_hint' is already applied with different arguments}} expected-note {{previous attribute is here}}
+
+// Show that the attribute works on member functions.
+class Functor_1 {
+public:
+  [[sycl::work_group_size_hint(16, 1, 1)]] [[sycl::work_group_size_hint(16, 1, 1)]] void operator()() const;
+  [[sycl::work_group_size_hint(16, 1, 1)]] [[sycl::work_group_size_hint(32, 1, 1)]] void operator()(int) const; // expected-warning {{attribute 'work_group_size_hint' is already applied with different arguments}} expected-note {{previous attribute is here}}
+};
+
+// Ensure that template arguments behave appropriately based on instantiations.
+template <int N>
+[[sycl::work_group_size_hint(N, 1, 1)]] void f8(); // #f8
+
+// Test that template redeclarations also get diagnosed properly.
+template <int X, int Y, int Z>
+[[sycl::work_group_size_hint(1, 1, 1)]] void f9(); // #f9prev
+
+template <int X, int Y, int Z>
+[[sycl::work_group_size_hint(X, Y, Z)]] void f9() {} // #f9
+
+// Test that a template redeclaration where the difference is known up front is
+// diagnosed immediately, even without instantiation.
+template <int X, int Y, int Z>
+[[sycl::work_group_size_hint(X, 1, Z)]] void f10(); // expected-note {{previous attribute is here}}
+template <int X, int Y, int Z>
+[[sycl::work_group_size_hint(X, 2, Z)]] void f10(); // expected-warning {{attribute 'work_group_size_hint' is already applied with different arguments}}
+
+#ifdef TRIGGER_ERROR
+[[sycl::work_group_size_hint(1, 2, 0)]] void f11(); // expected-error {{'work_group_size_hint' attribute requires a positive integral compile time constant expression}}
+#endif
+
+void instantiate() {
+  f8<1>(); // OK
+#ifdef TRIGGER_ERROR
+  // expected-error@#f8 {{'work_group_size_hint' attribute requires a positive integral compile time constant expression}}
+  f8<-1>(); // expected-note {{in instantiation}}
+  // expected-error@#f8 {{'work_group_size_hint' attribute requires a positive integral compile time constant expression}}
+  f8<0>(); // expected-note {{in instantiation}}
+#endif
+
+  f9<1, 1, 1>(); // OK, args are the same on the redecl.
+
+  // expected-warning@#f9 {{attribute 'work_group_size_hint' is already applied with different arguments}}
+  // expected-note@#f9prev {{previous attribute is here}}
+  f9<1, 2, 3>(); // expected-note {{in instantiation}}
+}
+
+// Show that the attribute works on member functions.
+class Functor16x2x1 {
+public:
+  [[sycl::work_group_size_hint(16, 2, 1)]] void operator()() const {};
+};
+
+// CHECK: CXXRecordDecl {{.*}} {{.*}}Functor16x2x1
+// CHECK: WorkGroupSizeHintAttr {{.*}}
+// CHECK-NEXT:  ConstantExpr{{.*}}'int'
+// CHECK-NEXT:  value: Int 16
+// CHECK-NEXT:  IntegerLiteral{{.*}}16{{$}}
+// CHECK-NEXT:  ConstantExpr{{.*}}'int'
+// CHECK-NEXT:  value: Int 2
+// CHECK-NEXT:  IntegerLiteral{{.*}}2{{$}}
+// CHECK-NEXT:  ConstantExpr{{.*}}'int'
+// CHECK-NEXT:  value: Int 1
+// CHECK-NEXT:  IntegerLiteral{{.*}}1{{$}}
+
+class Functor4x4x4 {
+public:
+  [[sycl::work_group_size_hint(4, 4, 4)]] void operator()() const {};
+};
+
+// Checking whether propagation of the attribute happens or not, according to the SYCL version.
+#if defined(EXPECT_PROP) // if attribute is propagated, then we expect errors here
+void f8x8x8(){};
+#else // otherwise no error
+[[sycl::work_group_size_hint(8, 8, 8)]] void f8x8x8(){};
+#endif
+class FunctorNoProp {
+public:
+  void operator()() const {
+    f8x8x8();
+  };
+};
+
+void invoke() {
+  Functor16x2x1 f16x2x1;
+  Functor4x4x4 f4x4x4;
+
+  sycl::queue q;
+
+  q.submit([&](sycl::handler &h) {
+    h.single_task<class kernel_1>(f16x2x1);
+    // CHECK: FunctionDecl {{.*}} {{.*}}kernel_1
+    // CHECK: WorkGroupSizeHintAttr {{.*}}
+    // CHECK-NEXT:  ConstantExpr{{.*}}'int'
+    // CHECK-NEXT:  value: Int 16
+    // CHECK-NEXT:  IntegerLiteral{{.*}}16{{$}}
+    // CHECK-NEXT:  ConstantExpr{{.*}}'int'
+    // CHECK-NEXT:  value: Int 2
+    // CHECK-NEXT:  IntegerLiteral{{.*}}2{{$}}
+    // CHECK-NEXT:  ConstantExpr{{.*}}'int'
+    // CHECK-NEXT:  value: Int 1
+    // CHECK-NEXT:  IntegerLiteral{{.*}}1{{$}}
+
+    // Checking that attributes are propagated to the kernel from functions in SYCL 1.2.1 mode.
+#ifdef EXPECT_PROP
+    h.single_task<class kernel_2>([=]() {
+      f4x4x4();
+    });
+#else 
+    // Otherwise using a functor that has the required attributes 
+    h.single_task<class kernel_2>(f4x4x4);
+#endif
+    // CHECK: FunctionDecl {{.*}} {{.*}}kernel_2
+    // CHECK: WorkGroupSizeHintAttr {{.*}}
+    // CHECK-NEXT:  ConstantExpr{{.*}}'int'
+    // CHECK-NEXT:  value: Int 4
+    // CHECK-NEXT:  IntegerLiteral{{.*}}4{{$}}
+    // CHECK-NEXT:  ConstantExpr{{.*}}'int'
+    // CHECK-NEXT:  value: Int 4
+    // CHECK-NEXT:  IntegerLiteral{{.*}}4{{$}}
+    // CHECK-NEXT:  ConstantExpr{{.*}}'int'
+    // CHECK-NEXT:  value: Int 4
+    // CHECK-NEXT:  IntegerLiteral{{.*}}4{{$}}
+
+    // Check that conflicts are reported if the attribute is propagated in SYCL 1.2.1 mode.
+
+    FunctorNoProp fNoProp;
+    h.single_task<class kernel_3>(fNoProp);
+    // CHECK: FunctionDecl {{.*}} {{.*}}kernel_3
+    // CHECK-NOT: WorkGroupSizeHintAttr
+
+  });
+
+  // FIXME: Add tests with the C++23 lambda attribute syntax.
+}

clang/test/SemaSYCL/check-work-group-size-hint-host.cpp

@@ -0,0 +1,13 @@
+// RUN: %clang_cc1 -fsycl-is-host -fsyntax-only -verify %s
+
+// Host checks for the work_group_size_hint attribute.
+class Functor16x2x1 {
+public:
+  [[sycl::work_group_size_hint(16, 2, 1)]] void operator()() const {};
+};
+
+// The GNU spelling is deprecated in SYCL mode, but otherwise these attributes
+// have the same semantics.
+[[sycl::work_group_size_hint(4, 1, 1)]] void f4x1x1();
+__attribute__((work_group_size_hint(4, 1, 1))) void f4x1x1(); // expected-warning {{attribute 'work_group_size_hint' is deprecated}} \
+                                                          // expected-note {{did you mean to use '[[sycl::work_group_size_hint]]' instead?}}