From 0bf5ea881154bba9ba9c5f830925fd7c49f6c379 Mon Sep 17 00:00:00 2001
From: Erich Keane <erich.keane@intel.com>
Date: Mon, 29 Mar 2021 09:38:28 -0700
Subject: [PATCH 01/22] Implemented SYCL 2020 sub-group size functionality.

As specified here:
https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/SubGroup/SYCL_INTEL_sub_group.asciidoc#attributes

This patch implements the named_sub_group_size attribute as well as the
command line parameter, and creates a new spelling of
reqd_sub_group_size (sub_group_size) to work like the SYCL 2020 version.
---
 clang/include/clang/Basic/Attr.td             |  24 +++-
 clang/include/clang/Basic/AttrDocs.td         |  44 +++++++
 .../clang/Basic/DiagnosticSemaKinds.td        |   4 +
 clang/include/clang/Basic/LangOptions.def     |   8 ++
 clang/include/clang/Basic/LangOptions.h       |   7 +
 clang/include/clang/Driver/Options.td         |   7 +
 clang/lib/CodeGen/CodeGenFunction.cpp         |  41 +++++-
 clang/lib/Frontend/CompilerInvocation.cpp     |  22 ++++
 clang/lib/Sema/SemaDeclAttr.cpp               |  31 +++++
 clang/lib/Sema/SemaSYCL.cpp                   | 122 +++++++++++++++++-
 ...a-attribute-supported-attributes-list.test |   1 +
 11 files changed, 301 insertions(+), 10 deletions(-)

diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td
index 600f94f19ab18..936c0f63f64cc 100644
--- a/clang/include/clang/Basic/Attr.td
+++ b/clang/include/clang/Basic/Attr.td
@@ -1415,13 +1415,33 @@ def LoopUnrollHint : StmtAttr {
 }
 
 def IntelReqdSubGroupSize: InheritableAttr {
-  let Spellings = [GNU<"intel_reqd_sub_group_size">,
-                   CXX11<"intel", "reqd_sub_group_size">];
+  let Spellings = [
+    GNU<"intel_reqd_sub_group_size">, CXX11<"intel", "reqd_sub_group_size">,
+    CXX11<"intel", "sub_group_size"> // SYCL2020 spelling.
+  ];
   let Args = [ExprArgument<"Value">];
   let Subjects = SubjectList<[Function], ErrorDiag>;
   let Documentation = [IntelReqdSubGroupSizeDocs];
   let LangOpts = [OpenCL, SYCLIsDevice, SilentlyIgnoreSYCLIsHost];
   let SupportsNonconformingLambdaSyntax = 1;
+  let AdditionalMembers = [{
+    bool isSYCL2020Spelling() const {
+      return getSemanticSpelling() == CXX11_intel_sub_group_size;
+    }
+  }];
+}
+
+def IntelNamedSubGroupSize : InheritableAttr {
+  let Spellings = [CXX11<"intel", "named_sub_group_size">];
+  let Args = [IdentifierArgument<"Type">];
+  let Subjects = SubjectList<[Function], ErrorDiag>;
+  let Documentation = [IntelNamedSubGroupSizeDocs];
+  let LangOpts = [OpenCL, SYCLIsDevice, SilentlyIgnoreSYCLIsHost];
+  let SupportsNonconformingLambdaSyntax = 1;
+  let AdditionalMembers = [{
+    bool isPrimary() const { return getType()->isStr("primary"); }
+    bool isAuto() const { return getType()->isStr("auto"); }
+}];
 }
 
 // This attribute is both a type attribute, and a declaration attribute (for
diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td
index 8a3b2e207b58d..3ebc4480185b9 100644
--- a/clang/include/clang/Basic/AttrDocs.td
+++ b/clang/include/clang/Basic/AttrDocs.td
@@ -4254,6 +4254,10 @@ The [[intel::reqd_sub_group_size(n)]] attribute indicates that the kernel must
 be compiled and executed with a sub-group of size n. The value of n must be set
 to a sub-group size supported by the device, or device compilation will fail.
 
+The ``[[intel::sub_group_size(n)]]`` attribute has the same affect as the other
+forms of this, except it follows the SYCL 2020 Attribute Rules. See the
+``[[intel::named_sub_group_size(NAME)]]`` documentation for clarification.
+
 In addition to device functions, the required sub-group size attribute may also
 be specified in the definition of a named functor object and lambda functions,
 as in the examples below:
@@ -4277,6 +4281,46 @@ See Sub-groups for NDRange Parallelism proposal in sycl/doc/extensions/sub_group
   }];
 }
 
+def IntelNamedSubGroupSizeDocs : Documentation {
+  let Category = DocCatFunction;
+  let Content = [{
+The ``[[intel::named_sub_group_size(NAME)]]`` attribute works with the
+``[[intel::sub_group_size(N)]]`` attribute to define the specific subgroup size
+for the kernel. The ``[[intel::named_sub_group_size(NAME)]]`` form accepts a
+required parameter of either ``auto`` or ``primary``.
+
+``auto`` specifies that the implementation is free to select any of the valid
+sub-group sizes associated with the device to which the kernel is submitted.
+``primary`` specifies that the implementation should select the device's primary
+sub-group size as reported by ``info::device::primary_sub_group_size``.
+
+This attribute may not be combined with ``[[intel::sub_group_size(N)]]``, as
+the two attributes have different meanings.
+
+In addition to the attributes, a default sub-group size strategy may be
+specified by the command line parameter ``-fsycl-default-sub-group-size`` flag,
+which accepts either ``auto``, ``primary``, or a default size as an integer.
+These values match and have the same behavior as the ``auto``, ``primary``, and
+``[[intel::sub_group_size(N)]]`` values respectively.
+
+SYCL 2020 Attribute Rules:
+SYCL 2020 specifies that kernel-type attributes should only be specified on the
+kernel or a ``SYCL_EXTERNAL`` function. This implementation permits these
+attributes to appear on all function declarations for the purposes of
+self-documenting declarations. However, these attributes must match the kernel's
+sub-group size as configured by the command line, or via an attribute
+specifically.
+
+In addition to the SYCL 2020 Attribute Rules, this attribute and the
+``[[intel::sub_group_size(N)]]`` attribute also require that any
+``SYCL_EXTERNAL`` functions defined in a different translation unit must have a
+matching sub-group size specification, so ``SYCL_EXTERNAL`` functions not
+defined in this translation unit must also have a matching sub-group
+specification to the kernel function that calls it.
+:xa
+}];
+}
+
 def OpenCLAccessDocs : Documentation {
   let Category = DocCatStmt;
   let Heading = "__read_only, __write_only, __read_write (read_only, write_only, read_write)";
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index bfafb3da5c30d..07e8e66de958e 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -11264,6 +11264,10 @@ def note_spelling_suggestion : Note<
   "did you mean to use %0 instead?">;
 def warn_attribute_requires_non_negative_integer_argument :
   Warning<warn_impcast_integer_sign.Text>, InGroup<AcceptedAttributes>;
+def err_sycl_mismatch_group_size
+    : Error<"%select{kernel called|undefined SYCL_EXTERNAL}0 function must "
+            "have a sub group size that matches size specified for the kernel">;
+def note_sycl_kernel_declared_here : Note<"kernel declared here">;
 
 // errors of expect.with.probability
 def err_probability_not_constant_float : Error<
diff --git a/clang/include/clang/Basic/LangOptions.def b/clang/include/clang/Basic/LangOptions.def
index 33942bbecf6ef..f45272634c17a 100644
--- a/clang/include/clang/Basic/LangOptions.def
+++ b/clang/include/clang/Basic/LangOptions.def
@@ -419,6 +419,14 @@ LANGOPT(RelativeCXXABIVTables, 1, 0,
 
 LANGOPT(ArmSveVectorBits, 32, 0, "SVE vector size in bits")
 
+ENUM_LANGOPT(DefaultSubGroupSizeType, SubGroupSizeType, 2,
+             SubGroupSizeType::None,
+             "Strategy via which sub-group is assigned for SYCL kernel "
+             "types if not overridden via attributes")
+
+VALUE_LANGOPT(DefaultSubGroupSize, 32, 0,
+              "If DefaultSubGroupSizeType is Integer contains the value")
+
 #undef LANGOPT
 #undef COMPATIBLE_LANGOPT
 #undef BENIGN_LANGOPT
diff --git a/clang/include/clang/Basic/LangOptions.h b/clang/include/clang/Basic/LangOptions.h
index f9b042338adb9..7088987557a97 100644
--- a/clang/include/clang/Basic/LangOptions.h
+++ b/clang/include/clang/Basic/LangOptions.h
@@ -263,6 +263,13 @@ class LangOptions : public LangOptionsBase {
     Single
   };
 
+  enum class SubGroupSizeType {
+    None,
+    Auto,
+    Primary,
+    Integer
+  };
+
 public:
   /// The used language standard.
   LangStandard::Kind LangStd;
diff --git a/clang/include/clang/Driver/Options.td b/clang/include/clang/Driver/Options.td
index 84cec02d99628..e1cab4547c484 100644
--- a/clang/include/clang/Driver/Options.td
+++ b/clang/include/clang/Driver/Options.td
@@ -5546,6 +5546,13 @@ def sycl_std_EQ : Joined<["-"], "sycl-std=">, Group<sycl_Group>,
   AutoNormalizeEnum,
   ShouldParseIf<!strconcat(fsycl_is_device.KeyPath, "||", fsycl_is_host.KeyPath)>;
 
+def fsycl_default_sub_group_size
+    : Separate<["-"], "fsycl-default-sub-group-size">,
+      HelpText<"Set the default sub group size for SYCL kernels">;
+def fsycl_default_sub_group_size_EQ
+    : Joined<["-"], "fsycl-default-sub-group-size=">,
+      Alias<fsycl_default_sub_group_size>;
+
 defm cuda_approx_transcendentals : BoolFOption<"cuda-approx-transcendentals",
   LangOpts<"CUDADeviceApproxTranscendentals">, DefaultFalse,
   PosFlag<SetTrue, [CC1Option], "Use">, NegFlag<SetFalse, [], "Don't use">,
diff --git a/clang/lib/CodeGen/CodeGenFunction.cpp b/clang/lib/CodeGen/CodeGenFunction.cpp
index 7343c96890e86..60c9ac3f9b747 100644
--- a/clang/lib/CodeGen/CodeGenFunction.cpp
+++ b/clang/lib/CodeGen/CodeGenFunction.cpp
@@ -640,15 +640,50 @@ void CodeGenFunction::EmitOpenCLKernelMetadata(const FunctionDecl *FD,
                     llvm::MDNode::get(Context, AttrMDArgs));
   }
 
-  if (const IntelReqdSubGroupSizeAttr *A =
-          FD->getAttr<IntelReqdSubGroupSizeAttr>()) {
-    const auto *CE = dyn_cast<ConstantExpr>(A->getValue());
+  bool IsKernelOrDevice =
+      FD->hasAttr<SYCLKernelAttr>() || FD->hasAttr<SYCLDeviceAttr>();
+  const IntelReqdSubGroupSizeAttr *ReqSubGroup =
+      FD->getAttr<IntelReqdSubGroupSizeAttr>();
+
+  // To support the SYCL 2020 spelling with no propogation, only emit for
+  // kernel-or-device when that spelling, fall-back to old behavior.
+  if (ReqSubGroup && (IsKernelOrDevice || !ReqSubGroup->isSYCL2020Spelling())) {
+    const auto *CE = dyn_cast<ConstantExpr>(ReqSubGroup->getValue());
     assert(CE && "Not an integer constant expression");
     Optional<llvm::APSInt> ArgVal = CE->getResultAsAPSInt();
     llvm::Metadata *AttrMDArgs[] = {llvm::ConstantAsMetadata::get(
         Builder.getInt32(ArgVal->getSExtValue()))};
     Fn->setMetadata("intel_reqd_sub_group_size",
                     llvm::MDNode::get(Context, AttrMDArgs));
+  } else if (IsKernelOrDevice &&
+             CGM.getLangOpts().getDefaultSubGroupSizeType() ==
+                 LangOptions::SubGroupSizeType::Integer) {
+    llvm::Metadata *AttrMDArgs[] = {llvm::ConstantAsMetadata::get(
+        Builder.getInt32(CGM.getLangOpts().DefaultSubGroupSize))};
+    Fn->setMetadata("intel_reqd_sub_group_size",
+                    llvm::MDNode::get(Context, AttrMDArgs));
+  }
+
+  // SCYL2020 doesn't propogate attributes, so don't put it in an intermedate
+  // location.
+  if (IsKernelOrDevice) {
+    if (const IntelNamedSubGroupSizeAttr *A =
+            FD->getAttr<IntelNamedSubGroupSizeAttr>()) {
+      llvm::Metadata *AttrMDArgs[] = {
+          llvm::MDString::get(Context, A->isPrimary() ? "primary" : "auto")};
+      Fn->setMetadata("intel_reqd_sub_group_size",
+                      llvm::MDNode::get(Context, AttrMDArgs));
+    } else if (CGM.getLangOpts().getDefaultSubGroupSizeType() ==
+               LangOptions::SubGroupSizeType::Auto) {
+      llvm::Metadata *AttrMDArgs[] = {llvm::MDString::get(Context, "auto")};
+      Fn->setMetadata("intel_reqd_sub_group_size",
+                      llvm::MDNode::get(Context, AttrMDArgs));
+    } else if (CGM.getLangOpts().getDefaultSubGroupSizeType() ==
+               LangOptions::SubGroupSizeType::Primary) {
+      llvm::Metadata *AttrMDArgs[] = {llvm::MDString::get(Context, "primary")};
+      Fn->setMetadata("intel_reqd_sub_group_size",
+                      llvm::MDNode::get(Context, AttrMDArgs));
+    }
   }
 
   if (FD->hasAttr<SYCLSimdAttr>()) {
diff --git a/clang/lib/Frontend/CompilerInvocation.cpp b/clang/lib/Frontend/CompilerInvocation.cpp
index 7337285688383..c8f9a6291af98 100644
--- a/clang/lib/Frontend/CompilerInvocation.cpp
+++ b/clang/lib/Frontend/CompilerInvocation.cpp
@@ -3728,6 +3728,28 @@ bool CompilerInvocation::ParseLangArgsImpl(LangOptions &Opts, ArgList &Args,
     }
   }
 
+  // Parse SYCL Default Sub group size.
+  if (const Arg *A = Args.getLastArg(OPT_fsycl_default_sub_group_size)) {
+    StringRef Value = A->getValue();
+    Opts.setDefaultSubGroupSizeType(
+        llvm::StringSwitch<LangOptions::SubGroupSizeType>(Value)
+            .Case("auto", LangOptions::SubGroupSizeType::Auto)
+            .Case("primary", LangOptions::SubGroupSizeType::Primary)
+            .Default(LangOptions::SubGroupSizeType::Integer));
+
+    if (Opts.getDefaultSubGroupSizeType() ==
+        LangOptions::SubGroupSizeType::Integer) {
+      int64_t IntResult;
+      if (Value.getAsInteger(10, IntResult)) {
+        Opts.DefaultSubGroupSize = IntResult;
+      } else {
+        Diags.Report(diag::err_drv_invalid_value)
+            << A->getAsString(Args) << A->getValue();
+        Opts.setDefaultSubGroupSizeType(LangOptions::SubGroupSizeType::None);
+      }
+    }
+  }
+
   Opts.DeclareSPIRVBuiltins = Args.hasArg(OPT_fdeclare_spirv_builtins);
 
   // These need to be parsed now. They are used to set OpenCL defaults.
diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index 3da76457434a3..f3a13f9801653 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -3210,6 +3210,11 @@ static void handleWorkGroupSizeHint(Sema &S, Decl *D, const ParsedAttr &AL) {
 
 void Sema::AddIntelReqdSubGroupSize(Decl *D, const AttributeCommonInfo &CI,
                                     Expr *E) {
+  if (checkAttrMutualExclusion<IntelNamedSubGroupSizeAttr>(*this, D, CI))
+    return;
+  if (checkAttrMutualExclusion<SYCLSimdAttr>(*this, D, CI))
+    return;
+
   if (!E->isValueDependent()) {
     // Validate that we have an integer constant expression and then store the
     // converted constant expression into the semantic attribute so that we
@@ -3269,6 +3274,29 @@ static void handleIntelReqdSubGroupSize(Sema &S, Decl *D,
   S.AddIntelReqdSubGroupSize(D, AL, E);
 }
 
+static void handleIntelNamedSubGroupSize(Sema &S, Decl *D,
+                                         const ParsedAttr &AL) {
+  if (checkAttrMutualExclusion<IntelReqdSubGroupSizeAttr>(S, D, AL))
+    return;
+  if (checkAttrMutualExclusion<SYCLSimdAttr>(S, D, AL))
+    return;
+
+  if (!AL.isArgIdent(0)) {
+    S.Diag(AL.getArgAsExpr(0)->getBeginLoc(),
+           diag::err_attribute_argument_type)
+        << AL << AANT_ArgumentIdentifier;
+    return;
+  }
+
+  IdentifierLoc *IL = AL.getArgAsIdent(0);
+  if (!IL->Ident->isStr("auto") && !IL->Ident->isStr("primary")) {
+    S.Diag(IL->Loc, diag::warn_attribute_type_not_supported) << AL << IL->Ident;
+    return;
+  }
+
+  D->addAttr(IntelNamedSubGroupSizeAttr::Create(S.Context, IL->Ident, AL));
+}
+
 void Sema::AddSYCLIntelNumSimdWorkItemsAttr(Decl *D,
                                             const AttributeCommonInfo &CI,
                                             Expr *E) {
@@ -9276,6 +9304,9 @@ static void ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D,
   case ParsedAttr::AT_IntelReqdSubGroupSize:
     handleIntelReqdSubGroupSize(S, D, AL);
     break;
+  case ParsedAttr::AT_IntelNamedSubGroupSize:
+    handleIntelNamedSubGroupSize(S, D, AL);
+    break;
   case ParsedAttr::AT_SYCLIntelNumSimdWorkItems:
     handleSYCLIntelNumSimdWorkItemsAttr(S, D, AL);
     break;
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 9ffa0ce900698..bbcbd6d0ad896 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -321,11 +321,12 @@ 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, ReqdWorkGroupSizeAttr,
-               SYCLIntelKernelArgsRestrictAttr, SYCLIntelNumSimdWorkItemsAttr,
-               SYCLIntelSchedulerTargetFmaxMhzAttr,
-               SYCLIntelMaxWorkGroupSizeAttr, SYCLIntelMaxGlobalWorkDimAttr,
-               SYCLIntelNoGlobalWorkOffsetAttr, SYCLSimdAttr>(A);
+        return isa<IntelReqdSubGroupSizeAttr, IntelNamedSubGroupSizeAttr,
+                   ReqdWorkGroupSizeAttr, SYCLIntelKernelArgsRestrictAttr,
+                   SYCLIntelNumSimdWorkItemsAttr,
+                   SYCLIntelSchedulerTargetFmaxMhzAttr,
+                   SYCLIntelMaxWorkGroupSizeAttr, SYCLIntelMaxGlobalWorkDimAttr,
+                   SYCLIntelNoGlobalWorkOffsetAttr, SYCLSimdAttr>(A);
   });
 
   // Allow the kernel attribute "use_stall_enable_clusters" only on lambda
@@ -3405,6 +3406,105 @@ void Sema::ConstructOpenCLKernel(FunctionDecl *KernelCallerFunc,
   }
 }
 
+// Figure out the sub-group for the this function.  First we check the
+// attributes, then the global settings.
+static std::pair<LangOptions::SubGroupSizeType, int64_t>
+CalcEffectiveSubGroup(ASTContext &Ctx, const LangOptions &LO,
+                      FunctionDecl *FD) {
+  if (const auto *A = FD->getAttr<IntelReqdSubGroupSizeAttr>()) {
+    int64_t Val = getIntExprValue(A->getValue(), Ctx);
+    return {LangOptions::SubGroupSizeType::Integer, Val};
+  }
+
+  if(const auto *A = FD->getAttr<IntelNamedSubGroupSizeAttr>()) {
+    if (A->isPrimary())
+      return {LangOptions::SubGroupSizeType::Primary, 0};
+    return {LangOptions::SubGroupSizeType::Auto, 0};
+  }
+
+  // Return the global settings.
+  return {LO.getDefaultSubGroupSizeType(),
+          static_cast<uint64_t>(LO.DefaultSubGroupSize)};
+}
+
+static SourceLocation GetSubGroupLoc(FunctionDecl *FD) {
+  if (const auto *A = FD->getAttr<IntelReqdSubGroupSizeAttr>())
+    return A->getLocation();
+  if (const auto *A = FD->getAttr<IntelNamedSubGroupSizeAttr>())
+    return A->getLocation();
+  return SourceLocation{};
+}
+
+static void CheckSYCL2020SubGroupSizes(Sema &S, FunctionDecl *SYCLKernel,
+                                       FunctionDecl *FD) {
+  // If they are the same, no error.
+  if (CalcEffectiveSubGroup(S.Context, S.getLangOpts(), SYCLKernel) ==
+      CalcEffectiveSubGroup(S.Context, S.getLangOpts(), FD))
+    return;
+
+  // Else we need to figure out why they don't match.
+  SourceLocation FDAttrLoc = GetSubGroupLoc(FD);
+  SourceLocation KernelAttrLoc = GetSubGroupLoc(SYCLKernel);
+
+  if (FDAttrLoc.isValid()) {
+    // This side was caused by an attribute.
+    S.Diag(FDAttrLoc, diag::err_sycl_mismatch_group_size)
+        << /*kernel called*/ 0;
+
+    if (KernelAttrLoc.isValid()) {
+      S.Diag(KernelAttrLoc, diag::note_conflicting_attribute);
+    } else {
+      // Kernel is 'default'.
+      S.Diag(SYCLKernel->getLocation(), diag::note_sycl_kernel_declared_here);
+    }
+    return;
+  }
+
+  // Else this doesn't have an attribute, which can only be caused by this being
+  // an undefined SYCL_EXTERNAL, and the kernel has an attribute that conflicts.
+  assert(KernelAttrLoc.isValid() && "Kernel doesn't have attribute either?");
+  S.Diag(FD->getLocation(), diag::err_sycl_mismatch_group_size)
+      << /*undefined SYCL_EXTERNAL*/ 1;
+  S.Diag(KernelAttrLoc, diag::note_conflicting_attribute);
+}
+
+// Check SYCL2020 Attributes.  2020 attributes don't propogate, they are only
+// valid if they match the attribute on the kernel. Note that this is a slight
+// difference from what the spec says, which says these attributes are only
+// valid on SYCL Kernels and SYCL_EXTERNAL, but we felt that for
+// self-documentation purposes that it would be nice to be able to repeat these
+// on subsequent functions.
+static void
+CheckSYCL2020Attributes(Sema &S, FunctionDecl *SYCLKernel,
+                        llvm::SmallPtrSetImpl<FunctionDecl *> &CalledFuncs) {
+
+  for (auto *FD : CalledFuncs) {
+    for (auto *Attr : FD->attrs()) {
+      switch (Attr->getKind()) {
+      case attr::Kind::IntelReqdSubGroupSize:
+        if (const auto *A = cast<IntelReqdSubGroupSizeAttr>(Attr))
+          // Pre SYCL2020 spellings handled during collection.
+          if (!A->isSYCL2020Spelling())
+            break;
+        LLVM_FALLTHROUGH;
+      case attr::Kind::IntelNamedSubGroupSize:
+        CheckSYCL2020SubGroupSizes(S, SYCLKernel, FD);
+        break;
+      case attr::Kind::SYCLDevice:
+        // If a SYCL_EXTERNAL function is not defined in this TU, its necessary
+        // that it has a compatible sub-group-size. Don't diagnose if it has a
+        // sub-group attribute, we can count on the other checks to catch this.
+        if (!FD->isDefined() && !FD->hasAttr<IntelReqdSubGroupSizeAttr>() &&
+            !FD->hasAttr<IntelNamedSubGroupSizeAttr>())
+          CheckSYCL2020SubGroupSizes(S, SYCLKernel, FD);
+        break;
+      default:
+        break;
+      }
+    }
+  }
+}
+
 void Sema::MarkDevice(void) {
   // Create the call graph so we can detect recursion and check the validity
   // of new operator overrides. Add the kernel function itself in case
@@ -3434,12 +3534,21 @@ void Sema::MarkDevice(void) {
       FunctionDecl *KernelBody =
           Marker.CollectPossibleKernelAttributes(SYCLKernel, Attrs);
 
+      CheckSYCL2020Attributes(*this, SYCLKernel, VisitedSet);
+
       for (auto *A : Attrs) {
         switch (A->getKind()) {
         case attr::Kind::IntelReqdSubGroupSize: {
           auto *Attr = cast<IntelReqdSubGroupSizeAttr>(A);
+
+          // SYCL2020 spelling, handled elsewhere.
+          if (Attr->isSYCL2020Spelling())
+            break;
+
           const auto *KBSimdAttr =
               KernelBody ? KernelBody->getAttr<SYCLSimdAttr>() : nullptr;
+          // If 'Existing' is a 2020 spelling, this should still conflict, so no
+          // special work is done here.
           if (auto *Existing =
                   SYCLKernel->getAttr<IntelReqdSubGroupSizeAttr>()) {
             if (getIntExprValue(Existing->getValue(), getASTContext()) !=
@@ -3531,6 +3640,9 @@ void Sema::MarkDevice(void) {
           break;
         }
         // TODO: vec_len_hint should be handled here
+        case attr::Kind::IntelNamedSubGroupSize:
+          // Nothing to do here, SYCL 2020 attr only.
+          break;
         default:
           // Seeing this means that CollectPossibleKernelAttributes was
           // updated while this switch wasn't...or something went wrong
diff --git a/clang/test/Misc/pragma-attribute-supported-attributes-list.test b/clang/test/Misc/pragma-attribute-supported-attributes-list.test
index ad21694ff4d4c..3b29397a855b8 100644
--- a/clang/test/Misc/pragma-attribute-supported-attributes-list.test
+++ b/clang/test/Misc/pragma-attribute-supported-attributes-list.test
@@ -71,6 +71,7 @@
 // CHECK-NEXT: IBAction (SubjectMatchRule_objc_method_is_instance)
 // CHECK-NEXT: IFunc (SubjectMatchRule_function)
 // CHECK-NEXT: InitPriority (SubjectMatchRule_variable)
+// CHECK-NEXT: IntelNamedSubGroupSize (SubjectMatchRule_function)
 // CHECK-NEXT: IntelReqdSubGroupSize (SubjectMatchRule_function)
 // CHECK-NEXT: InternalLinkage (SubjectMatchRule_variable, SubjectMatchRule_function, SubjectMatchRule_record)
 // CHECK-NEXT: LTOVisibilityPublic (SubjectMatchRule_record)

From 25a28b3022bebc32d1be30484097434603a9dff9 Mon Sep 17 00:00:00 2001
From: Erich Keane <erich.keane@intel.com>
Date: Tue, 30 Mar 2021 09:02:47 -0700
Subject: [PATCH 02/22] Apply clang-format changes

---
 clang/include/clang/Basic/LangOptions.h |  7 +------
 clang/lib/Sema/SemaDeclAttr.cpp         |  3 +--
 clang/lib/Sema/SemaSYCL.cpp             | 12 ++++++------
 3 files changed, 8 insertions(+), 14 deletions(-)

diff --git a/clang/include/clang/Basic/LangOptions.h b/clang/include/clang/Basic/LangOptions.h
index 7088987557a97..2c8b45cebbe78 100644
--- a/clang/include/clang/Basic/LangOptions.h
+++ b/clang/include/clang/Basic/LangOptions.h
@@ -263,12 +263,7 @@ class LangOptions : public LangOptionsBase {
     Single
   };
 
-  enum class SubGroupSizeType {
-    None,
-    Auto,
-    Primary,
-    Integer
-  };
+  enum class SubGroupSizeType { None, Auto, Primary, Integer };
 
 public:
   /// The used language standard.
diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index f3a13f9801653..1904207a46c6b 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -3282,8 +3282,7 @@ static void handleIntelNamedSubGroupSize(Sema &S, Decl *D,
     return;
 
   if (!AL.isArgIdent(0)) {
-    S.Diag(AL.getArgAsExpr(0)->getBeginLoc(),
-           diag::err_attribute_argument_type)
+    S.Diag(AL.getArgAsExpr(0)->getBeginLoc(), diag::err_attribute_argument_type)
         << AL << AANT_ArgumentIdentifier;
     return;
   }
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index bbcbd6d0ad896..8ec9e821de929 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -321,12 +321,12 @@ 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,
-                   SYCLIntelSchedulerTargetFmaxMhzAttr,
-                   SYCLIntelMaxWorkGroupSizeAttr, SYCLIntelMaxGlobalWorkDimAttr,
-                   SYCLIntelNoGlobalWorkOffsetAttr, SYCLSimdAttr>(A);
+    return isa<IntelReqdSubGroupSizeAttr, IntelNamedSubGroupSizeAttr,
+               ReqdWorkGroupSizeAttr, SYCLIntelKernelArgsRestrictAttr,
+               SYCLIntelNumSimdWorkItemsAttr,
+               SYCLIntelSchedulerTargetFmaxMhzAttr,
+               SYCLIntelMaxWorkGroupSizeAttr, SYCLIntelMaxGlobalWorkDimAttr,
+               SYCLIntelNoGlobalWorkOffsetAttr, SYCLSimdAttr>(A);
   });
 
   // Allow the kernel attribute "use_stall_enable_clusters" only on lambda

From 646a2270254ec64023c6fb76cc7026e731689ecf Mon Sep 17 00:00:00 2001
From: Erich Keane <erich.keane@intel.com>
Date: Tue, 30 Mar 2021 09:07:36 -0700
Subject: [PATCH 03/22] Another clang-format change?

---
 clang/lib/Sema/SemaSYCL.cpp | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 8ec9e821de929..93c53308fc418 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -3416,7 +3416,7 @@ CalcEffectiveSubGroup(ASTContext &Ctx, const LangOptions &LO,
     return {LangOptions::SubGroupSizeType::Integer, Val};
   }
 
-  if(const auto *A = FD->getAttr<IntelNamedSubGroupSizeAttr>()) {
+  if (const auto *A = FD->getAttr<IntelNamedSubGroupSizeAttr>()) {
     if (A->isPrimary())
       return {LangOptions::SubGroupSizeType::Primary, 0};
     return {LangOptions::SubGroupSizeType::Auto, 0};

From 5c03d1a921a24187cf849b8e3e8a2d3a8a5fda42 Mon Sep 17 00:00:00 2001
From: Erich Keane <erich.keane@intel.com>
Date: Tue, 30 Mar 2021 10:13:13 -0700
Subject: [PATCH 04/22] Fix all of Aaron's comments, plus change auto to
 automatic

The feature architect agreed to change auto to automatic due to
implementation issues.  Also changed over to EnumArgument.
---
 clang/include/clang/Basic/Attr.td             |  7 ++--
 clang/include/clang/Basic/AttrDocs.td         | 34 ++++++++++---------
 .../clang/Basic/DiagnosticSemaKinds.td        |  5 +--
 clang/include/clang/Sema/Sema.h               |  2 ++
 clang/lib/CodeGen/CodeGenFunction.cpp         | 13 ++++---
 clang/lib/Frontend/CompilerInvocation.cpp     |  2 +-
 clang/lib/Sema/SemaDecl.cpp                   |  2 ++
 clang/lib/Sema/SemaDeclAttr.cpp               | 32 +++++++++++++++--
 clang/lib/Sema/SemaSYCL.cpp                   |  2 +-
 9 files changed, 67 insertions(+), 32 deletions(-)

diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td
index 936c0f63f64cc..352102075569e 100644
--- a/clang/include/clang/Basic/Attr.td
+++ b/clang/include/clang/Basic/Attr.td
@@ -1433,15 +1433,12 @@ def IntelReqdSubGroupSize: InheritableAttr {
 
 def IntelNamedSubGroupSize : InheritableAttr {
   let Spellings = [CXX11<"intel", "named_sub_group_size">];
-  let Args = [IdentifierArgument<"Type">];
+  let Args = [EnumArgument<"Type", "SubGroupSizeType", ["automatic", "primary"],
+                           ["Auto", "Primary"]>];
   let Subjects = SubjectList<[Function], ErrorDiag>;
   let Documentation = [IntelNamedSubGroupSizeDocs];
   let LangOpts = [OpenCL, SYCLIsDevice, SilentlyIgnoreSYCLIsHost];
   let SupportsNonconformingLambdaSyntax = 1;
-  let AdditionalMembers = [{
-    bool isPrimary() const { return getType()->isStr("primary"); }
-    bool isAuto() const { return getType()->isStr("auto"); }
-}];
 }
 
 // This attribute is both a type attribute, and a declaration attribute (for
diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td
index 3ebc4480185b9..35c1fefba8a3c 100644
--- a/clang/include/clang/Basic/AttrDocs.td
+++ b/clang/include/clang/Basic/AttrDocs.td
@@ -4254,9 +4254,11 @@ The [[intel::reqd_sub_group_size(n)]] attribute indicates that the kernel must
 be compiled and executed with a sub-group of size n. The value of n must be set
 to a sub-group size supported by the device, or device compilation will fail.
 
-The ``[[intel::sub_group_size(n)]]`` attribute has the same affect as the other
-forms of this, except it follows the SYCL 2020 Attribute Rules. See the
-``[[intel::named_sub_group_size(NAME)]]`` documentation for clarification.
+The ``[[intel::sub_group_size(n)]]`` attribute has the same effect as the other
+attribute spellings, except that it follows the SYCL 2020 Attribute Rules. See
+ the ``[[intel::named_sub_group_size(NAME)]]`` documentation for clarification.
+
+This attribute is mutually exclusive with ``[[intel::named_sub_group_size(NAME)]]``.
 
 In addition to device functions, the required sub-group size attribute may also
 be specified in the definition of a named functor object and lambda functions,
@@ -4284,24 +4286,25 @@ See Sub-groups for NDRange Parallelism proposal in sycl/doc/extensions/sub_group
 def IntelNamedSubGroupSizeDocs : Documentation {
   let Category = DocCatFunction;
   let Content = [{
-The ``[[intel::named_sub_group_size(NAME)]]`` attribute works with the
-``[[intel::sub_group_size(N)]]`` attribute to define the specific subgroup size
-for the kernel. The ``[[intel::named_sub_group_size(NAME)]]`` form accepts a
-required parameter of either ``auto`` or ``primary``.
+The ``[[intel::named_sub_group_size(NAME)]]`` attribute works similar to
+``[[intel::sub_group_size(N)]]`` attribute in that it defines the specific
+subgroup size for the kernel. The ``[[intel::named_sub_group_size(NAME)]]`` form
+accepts a required parameter of either ``automatic`` or ``primary``.
 
-``auto`` specifies that the implementation is free to select any of the valid
-sub-group sizes associated with the device to which the kernel is submitted.
-``primary`` specifies that the implementation should select the device's primary
-sub-group size as reported by ``info::device::primary_sub_group_size``.
+``automatic`` specifies that the implementation is free to select any of the
+valid sub-group sizes associated with the device to which the kernel is
+submitted.  ``primary`` specifies that the implementation should select the
+device's primary sub-group size as reported by
+``info::device::primary_sub_group_size``.
 
 This attribute may not be combined with ``[[intel::sub_group_size(N)]]``, as
 the two attributes have different meanings.
 
 In addition to the attributes, a default sub-group size strategy may be
-specified by the command line parameter ``-fsycl-default-sub-group-size`` flag,
-which accepts either ``auto``, ``primary``, or a default size as an integer.
-These values match and have the same behavior as the ``auto``, ``primary``, and
-``[[intel::sub_group_size(N)]]`` values respectively.
+specified by the ``-fsycl-default-sub-group-size`` command line option, which
+accepts either ``automatic``, ``primary``, or a default size as an integer.
+These values match and have the same behavior as the ``automatic``, ``primary``,
+and ``[[intel::sub_group_size(N)]]`` values respectively.
 
 SYCL 2020 Attribute Rules:
 SYCL 2020 specifies that kernel-type attributes should only be specified on the
@@ -4317,7 +4320,6 @@ In addition to the SYCL 2020 Attribute Rules, this attribute and the
 matching sub-group size specification, so ``SYCL_EXTERNAL`` functions not
 defined in this translation unit must also have a matching sub-group
 specification to the kernel function that calls it.
-:xa
 }];
 }
 
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index 07e8e66de958e..b79ca565bebbc 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -11265,8 +11265,9 @@ def note_spelling_suggestion : Note<
 def warn_attribute_requires_non_negative_integer_argument :
   Warning<warn_impcast_integer_sign.Text>, InGroup<AcceptedAttributes>;
 def err_sycl_mismatch_group_size
-    : Error<"%select{kernel called|undefined SYCL_EXTERNAL}0 function must "
-            "have a sub group size that matches size specified for the kernel">;
+    : Error<"%select{kernel-called|undefined 'SYCL_EXTERNAL'}0 function must "
+            "have a sub group size that matches the size specified for the "
+            "kernel">;
 def note_sycl_kernel_declared_here : Note<"kernel declared here">;
 
 // errors of expect.with.probability
diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h
index 54cc16faf56e4..1a7b6d176582b 100644
--- a/clang/include/clang/Sema/Sema.h
+++ b/clang/include/clang/Sema/Sema.h
@@ -10211,6 +10211,8 @@ class Sema final {
                                 Expr *E);
   IntelReqdSubGroupSizeAttr *
   MergeIntelReqdSubGroupSizeAttr(Decl *D, const IntelReqdSubGroupSizeAttr &A);
+  IntelNamedSubGroupSizeAttr *
+  MergeIntelNamedSubGroupSizeAttr(Decl *D, const IntelNamedSubGroupSizeAttr &A);
   void AddSYCLIntelNumSimdWorkItemsAttr(Decl *D, const AttributeCommonInfo &CI,
                                         Expr *E);
   SYCLIntelNumSimdWorkItemsAttr *
diff --git a/clang/lib/CodeGen/CodeGenFunction.cpp b/clang/lib/CodeGen/CodeGenFunction.cpp
index 60c9ac3f9b747..bcbeaf9da7e08 100644
--- a/clang/lib/CodeGen/CodeGenFunction.cpp
+++ b/clang/lib/CodeGen/CodeGenFunction.cpp
@@ -645,7 +645,7 @@ void CodeGenFunction::EmitOpenCLKernelMetadata(const FunctionDecl *FD,
   const IntelReqdSubGroupSizeAttr *ReqSubGroup =
       FD->getAttr<IntelReqdSubGroupSizeAttr>();
 
-  // To support the SYCL 2020 spelling with no propogation, only emit for
+  // To support the SYCL 2020 spelling with no propagation, only emit for
   // kernel-or-device when that spelling, fall-back to old behavior.
   if (ReqSubGroup && (IsKernelOrDevice || !ReqSubGroup->isSYCL2020Spelling())) {
     const auto *CE = dyn_cast<ConstantExpr>(ReqSubGroup->getValue());
@@ -664,18 +664,21 @@ void CodeGenFunction::EmitOpenCLKernelMetadata(const FunctionDecl *FD,
                     llvm::MDNode::get(Context, AttrMDArgs));
   }
 
-  // SCYL2020 doesn't propogate attributes, so don't put it in an intermedate
+  // SCYL2020 doesn't propagate attributes, so don't put it in an intermediate
   // location.
   if (IsKernelOrDevice) {
     if (const IntelNamedSubGroupSizeAttr *A =
             FD->getAttr<IntelNamedSubGroupSizeAttr>()) {
-      llvm::Metadata *AttrMDArgs[] = {
-          llvm::MDString::get(Context, A->isPrimary() ? "primary" : "auto")};
+      llvm::Metadata *AttrMDArgs[] = {llvm::MDString::get(
+          Context, A->getType() == IntelNamedSubGroupSizeAttr::Primary
+                       ? "primary"
+                       : "automatic")};
       Fn->setMetadata("intel_reqd_sub_group_size",
                       llvm::MDNode::get(Context, AttrMDArgs));
     } else if (CGM.getLangOpts().getDefaultSubGroupSizeType() ==
                LangOptions::SubGroupSizeType::Auto) {
-      llvm::Metadata *AttrMDArgs[] = {llvm::MDString::get(Context, "auto")};
+      llvm::Metadata *AttrMDArgs[] = {
+          llvm::MDString::get(Context, "automatic")};
       Fn->setMetadata("intel_reqd_sub_group_size",
                       llvm::MDNode::get(Context, AttrMDArgs));
     } else if (CGM.getLangOpts().getDefaultSubGroupSizeType() ==
diff --git a/clang/lib/Frontend/CompilerInvocation.cpp b/clang/lib/Frontend/CompilerInvocation.cpp
index c8f9a6291af98..b79b06686ca2b 100644
--- a/clang/lib/Frontend/CompilerInvocation.cpp
+++ b/clang/lib/Frontend/CompilerInvocation.cpp
@@ -3733,7 +3733,7 @@ bool CompilerInvocation::ParseLangArgsImpl(LangOptions &Opts, ArgList &Args,
     StringRef Value = A->getValue();
     Opts.setDefaultSubGroupSizeType(
         llvm::StringSwitch<LangOptions::SubGroupSizeType>(Value)
-            .Case("auto", LangOptions::SubGroupSizeType::Auto)
+            .Case("automatic", LangOptions::SubGroupSizeType::Auto)
             .Case("primary", LangOptions::SubGroupSizeType::Primary)
             .Default(LangOptions::SubGroupSizeType::Integer));
 
diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp
index f64866ba96048..72c68c41aaea5 100644
--- a/clang/lib/Sema/SemaDecl.cpp
+++ b/clang/lib/Sema/SemaDecl.cpp
@@ -2620,6 +2620,8 @@ static bool mergeDeclAttribute(Sema &S, NamedDecl *D,
     NewAttr = S.mergeEnforceTCBLeafAttr(D, *TCBLA);
   else if (const auto *A = dyn_cast<IntelReqdSubGroupSizeAttr>(Attr))
     NewAttr = S.MergeIntelReqdSubGroupSizeAttr(D, *A);
+  else if (const auto *A = dyn_cast<IntelNamedSubGroupSizeAttr>(Attr))
+    NewAttr = S.MergeIntelNamedSubGroupSizeAttr(D, *A);
   else if (const auto *A = dyn_cast<SYCLIntelNumSimdWorkItemsAttr>(Attr))
     NewAttr = S.MergeSYCLIntelNumSimdWorkItemsAttr(D, *A);
   else if (const auto *A = dyn_cast<SYCLIntelSchedulerTargetFmaxMhzAttr>(Attr))
diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index 1904207a46c6b..e12e2d40eab58 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -3253,6 +3253,11 @@ void Sema::AddIntelReqdSubGroupSize(Decl *D, const AttributeCommonInfo &CI,
 IntelReqdSubGroupSizeAttr *
 Sema::MergeIntelReqdSubGroupSizeAttr(Decl *D,
                                      const IntelReqdSubGroupSizeAttr &A) {
+  if (checkAttrMutualExclusion<IntelNamedSubGroupSizeAttr>(*this, D, A))
+    return nullptr;
+  if (checkAttrMutualExclusion<SYCLSimdAttr>(*this, D, A))
+    return nullptr;
+
   // Check to see if there's a duplicate attribute with different values
   // already applied to the declaration.
   if (const auto *DeclAttr = D->getAttr<IntelReqdSubGroupSizeAttr>()) {
@@ -3274,6 +3279,27 @@ static void handleIntelReqdSubGroupSize(Sema &S, Decl *D,
   S.AddIntelReqdSubGroupSize(D, AL, E);
 }
 
+IntelNamedSubGroupSizeAttr *
+Sema::MergeIntelNamedSubGroupSizeAttr(Decl *D,
+                                      const IntelNamedSubGroupSizeAttr &A) {
+  if (checkAttrMutualExclusion<IntelNamedSubGroupSizeAttr>(*this, D, A))
+    return nullptr;
+  if (checkAttrMutualExclusion<SYCLSimdAttr>(*this, D, A))
+    return nullptr;
+
+  // Check to see if there's a duplicate attribute with different values
+  // already applied to the declaration.
+  if (const auto *DeclAttr = D->getAttr<IntelNamedSubGroupSizeAttr>()) {
+    if (DeclAttr->getType() != A.getType()) {
+      Diag(DeclAttr->getLoc(), diag::warn_duplicate_attribute) << &A;
+      Diag(A.getLoc(), diag::note_previous_attribute);
+      return nullptr;
+    }
+  }
+
+  return IntelNamedSubGroupSizeAttr::Create(Context, A.getType(), A);
+}
+
 static void handleIntelNamedSubGroupSize(Sema &S, Decl *D,
                                          const ParsedAttr &AL) {
   if (checkAttrMutualExclusion<IntelReqdSubGroupSizeAttr>(S, D, AL))
@@ -3287,13 +3313,15 @@ static void handleIntelNamedSubGroupSize(Sema &S, Decl *D,
     return;
   }
 
+  IntelNamedSubGroupSizeAttr::SubGroupSizeType SizeType;
   IdentifierLoc *IL = AL.getArgAsIdent(0);
-  if (!IL->Ident->isStr("auto") && !IL->Ident->isStr("primary")) {
+  if (!IntelNamedSubGroupSizeAttr::ConvertStrToSubGroupSizeType(
+          IL->Ident->getName(), SizeType)) {
     S.Diag(IL->Loc, diag::warn_attribute_type_not_supported) << AL << IL->Ident;
     return;
   }
 
-  D->addAttr(IntelNamedSubGroupSizeAttr::Create(S.Context, IL->Ident, AL));
+  D->addAttr(IntelNamedSubGroupSizeAttr::Create(S.Context, SizeType, AL));
 }
 
 void Sema::AddSYCLIntelNumSimdWorkItemsAttr(Decl *D,
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 93c53308fc418..5337ab3611ad3 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -3417,7 +3417,7 @@ CalcEffectiveSubGroup(ASTContext &Ctx, const LangOptions &LO,
   }
 
   if (const auto *A = FD->getAttr<IntelNamedSubGroupSizeAttr>()) {
-    if (A->isPrimary())
+    if (A->getType() == IntelNamedSubGroupSizeAttr::Primary)
       return {LangOptions::SubGroupSizeType::Primary, 0};
     return {LangOptions::SubGroupSizeType::Auto, 0};
   }

From 149a8326dbed3104b0993eafcb95a54c617494d8 Mon Sep 17 00:00:00 2001
From: Erich Keane <erich.keane@intel.com>
Date: Tue, 30 Mar 2021 11:11:55 -0700
Subject: [PATCH 05/22] Fix attr.td param spelling + nullptr placement based on
 review comments

---
 clang/include/clang/Basic/Attr.td | 2 +-
 clang/lib/Sema/SemaDeclAttr.cpp   | 4 ++--
 2 files changed, 3 insertions(+), 3 deletions(-)

diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td
index 352102075569e..c81180b01a6d5 100644
--- a/clang/include/clang/Basic/Attr.td
+++ b/clang/include/clang/Basic/Attr.td
@@ -1434,7 +1434,7 @@ def IntelReqdSubGroupSize: InheritableAttr {
 def IntelNamedSubGroupSize : InheritableAttr {
   let Spellings = [CXX11<"intel", "named_sub_group_size">];
   let Args = [EnumArgument<"Type", "SubGroupSizeType", ["automatic", "primary"],
-                           ["Auto", "Primary"]>];
+                           ["Automatic", "Primary"]>];
   let Subjects = SubjectList<[Function], ErrorDiag>;
   let Documentation = [IntelNamedSubGroupSizeDocs];
   let LangOpts = [OpenCL, SYCLIsDevice, SilentlyIgnoreSYCLIsHost];
diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index e12e2d40eab58..70f8662637767 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -3282,7 +3282,7 @@ static void handleIntelReqdSubGroupSize(Sema &S, Decl *D,
 IntelNamedSubGroupSizeAttr *
 Sema::MergeIntelNamedSubGroupSizeAttr(Decl *D,
                                       const IntelNamedSubGroupSizeAttr &A) {
-  if (checkAttrMutualExclusion<IntelNamedSubGroupSizeAttr>(*this, D, A))
+  if (checkAttrMutualExclusion<IntelReqdSubGroupSizeAttr>(*this, D, A))
     return nullptr;
   if (checkAttrMutualExclusion<SYCLSimdAttr>(*this, D, A))
     return nullptr;
@@ -3293,8 +3293,8 @@ Sema::MergeIntelNamedSubGroupSizeAttr(Decl *D,
     if (DeclAttr->getType() != A.getType()) {
       Diag(DeclAttr->getLoc(), diag::warn_duplicate_attribute) << &A;
       Diag(A.getLoc(), diag::note_previous_attribute);
-      return nullptr;
     }
+    return nullptr;
   }
 
   return IntelNamedSubGroupSizeAttr::Create(Context, A.getType(), A);

From 2c35a955b3688da6aef07c9671335758d1d0a199 Mon Sep 17 00:00:00 2001
From: Erich Keane <erich.keane@intel.com>
Date: Tue, 30 Mar 2021 11:12:26 -0700
Subject: [PATCH 06/22] Add test for semadeclattr checks, semasycl checks to
 come

---
 clang/test/SemaSYCL/sub-group-size.cpp | 44 ++++++++++++++++++++++++++
 1 file changed, 44 insertions(+)
 create mode 100644 clang/test/SemaSYCL/sub-group-size.cpp

diff --git a/clang/test/SemaSYCL/sub-group-size.cpp b/clang/test/SemaSYCL/sub-group-size.cpp
new file mode 100644
index 0000000000000..0fad692581994
--- /dev/null
+++ b/clang/test/SemaSYCL/sub-group-size.cpp
@@ -0,0 +1,44 @@
+// RUN: %clang_cc1 -internal-isystem %S/Inputs -fsycl-is-device -internal-isystem %S/Inputs -fsyntax-only -verify %s
+
+#include "Inputs/sycl.hpp"
+
+// expected-error@+2 {{'named_sub_group_size' and 'sub_group_size' attributes are not compatible}}
+// expected-note@+1 {{conflicting attribute is here}}
+[[intel::sub_group_size(1)]][[intel::named_sub_group_size(automatic)]]
+void f1();
+// expected-error@+2 {{'sub_group_size' and 'named_sub_group_size' attributes are not compatible}}
+// expected-note@+1 {{conflicting attribute is here}}
+[[intel::named_sub_group_size(primary)]][[intel::sub_group_size(1)]]
+void f2();
+
+// expected-error@+1 {{'sub_group_size' and 'named_sub_group_size' attributes are not compatible}}
+[[intel::sub_group_size(1)]]
+void f3();
+// expected-note@+1 {{conflicting attribute is here}}
+[[intel::named_sub_group_size(primary)]]
+void f3();
+
+// expected-error@+1 {{'named_sub_group_size' and 'sub_group_size' attributes are not compatible}}
+[[intel::named_sub_group_size(primary)]]
+void f4();
+// expected-note@+1 {{conflicting attribute is here}}
+[[intel::sub_group_size(1)]]
+void f4();
+
+// expected-note@+1 {{previous attribute is here}}
+[[intel::named_sub_group_size(automatic)]]
+void f5();
+
+// expected-warning@+1 {{attribute 'named_sub_group_size' is already applied with different arguments}}
+[[intel::named_sub_group_size(primary)]]
+void f5();
+
+[[intel::named_sub_group_size(automatic)]]
+void f6();
+
+[[intel::named_sub_group_size(automatic)]]
+void f6();
+
+// expected-warning@+1 {{'named_sub_group_size' attribute argument not supported: 'invalid'}}
+[[intel::named_sub_group_size(invalid)]]
+void f7();

From 6be3cccfce0bcc6f47125c34f15039b67fd41b94 Mon Sep 17 00:00:00 2001
From: Erich Keane <erich.keane@intel.com>
Date: Tue, 30 Mar 2021 12:33:50 -0700
Subject: [PATCH 07/22] A few more comments from Aaron, plus getting the
 SemaSYCL diagnostics to work, plus some additional tests

---
 clang/include/clang/Basic/Attr.td      |  7 +--
 clang/include/clang/Sema/Sema.h        |  1 +
 clang/lib/Sema/SemaDecl.cpp            |  2 +
 clang/lib/Sema/SemaDeclAttr.cpp        | 22 ++++++++-
 clang/lib/Sema/SemaSYCL.cpp            |  8 +++-
 clang/test/SemaSYCL/sub-group-size.cpp | 65 +++++++++++++++++++++++++-
 6 files changed, 96 insertions(+), 9 deletions(-)

diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td
index c81180b01a6d5..676b1758dcf93 100644
--- a/clang/include/clang/Basic/Attr.td
+++ b/clang/include/clang/Basic/Attr.td
@@ -1424,11 +1424,8 @@ def IntelReqdSubGroupSize: InheritableAttr {
   let Documentation = [IntelReqdSubGroupSizeDocs];
   let LangOpts = [OpenCL, SYCLIsDevice, SilentlyIgnoreSYCLIsHost];
   let SupportsNonconformingLambdaSyntax = 1;
-  let AdditionalMembers = [{
-    bool isSYCL2020Spelling() const {
-      return getSemanticSpelling() == CXX11_intel_sub_group_size;
-    }
-  }];
+  let Accessors =
+      [Accessor<"isSYCL2020Spelling", [CXX11<"intel", "sub_group_size">]>];
 }
 
 def IntelNamedSubGroupSize : InheritableAttr {
diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h
index 1a7b6d176582b..a23ecfe91533d 100644
--- a/clang/include/clang/Sema/Sema.h
+++ b/clang/include/clang/Sema/Sema.h
@@ -10213,6 +10213,7 @@ class Sema final {
   MergeIntelReqdSubGroupSizeAttr(Decl *D, const IntelReqdSubGroupSizeAttr &A);
   IntelNamedSubGroupSizeAttr *
   MergeIntelNamedSubGroupSizeAttr(Decl *D, const IntelNamedSubGroupSizeAttr &A);
+  SYCLSimdAttr *MergeSYCLSimdAttr(Decl *D, const SYCLSimdAttr &A);
   void AddSYCLIntelNumSimdWorkItemsAttr(Decl *D, const AttributeCommonInfo &CI,
                                         Expr *E);
   SYCLIntelNumSimdWorkItemsAttr *
diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp
index 72c68c41aaea5..47a76171125ac 100644
--- a/clang/lib/Sema/SemaDecl.cpp
+++ b/clang/lib/Sema/SemaDecl.cpp
@@ -2622,6 +2622,8 @@ static bool mergeDeclAttribute(Sema &S, NamedDecl *D,
     NewAttr = S.MergeIntelReqdSubGroupSizeAttr(D, *A);
   else if (const auto *A = dyn_cast<IntelNamedSubGroupSizeAttr>(Attr))
     NewAttr = S.MergeIntelNamedSubGroupSizeAttr(D, *A);
+  else if (const auto *A = dyn_cast<SYCLSimdAttr>(Attr))
+    NewAttr = S.MergeSYCLSimdAttr(D, *A);
   else if (const auto *A = dyn_cast<SYCLIntelNumSimdWorkItemsAttr>(Attr))
     NewAttr = S.MergeSYCLIntelNumSimdWorkItemsAttr(D, *A);
   else if (const auto *A = dyn_cast<SYCLIntelSchedulerTargetFmaxMhzAttr>(Attr))
diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index 70f8662637767..7d2145c10af6c 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -3279,6 +3279,7 @@ static void handleIntelReqdSubGroupSize(Sema &S, Decl *D,
   S.AddIntelReqdSubGroupSize(D, AL, E);
 }
 
+
 IntelNamedSubGroupSizeAttr *
 Sema::MergeIntelNamedSubGroupSizeAttr(Decl *D,
                                       const IntelNamedSubGroupSizeAttr &A) {
@@ -3324,6 +3325,25 @@ static void handleIntelNamedSubGroupSize(Sema &S, Decl *D,
   D->addAttr(IntelNamedSubGroupSizeAttr::Create(S.Context, SizeType, AL));
 }
 
+SYCLSimdAttr *Sema::MergeSYCLSimdAttr(Decl *D, const SYCLSimdAttr &A) {
+  if (checkAttrMutualExclusion<IntelReqdSubGroupSizeAttr>(*this, D, A))
+    return nullptr;
+  if (checkAttrMutualExclusion<IntelNamedSubGroupSizeAttr>(*this, D, A))
+    return nullptr;
+
+  return A.clone(Context);
+}
+
+static void handleSYCLSimdAttr(Sema &S, Decl *D,
+                                         const ParsedAttr &AL) {
+  if (checkAttrMutualExclusion<IntelReqdSubGroupSizeAttr>(S, D, AL))
+    return;
+  if (checkAttrMutualExclusion<IntelNamedSubGroupSizeAttr>(S, D, AL))
+    return;
+
+  handleSimpleAttribute<SYCLSimdAttr>(S, D, AL);
+}
+
 void Sema::AddSYCLIntelNumSimdWorkItemsAttr(Decl *D,
                                             const AttributeCommonInfo &CI,
                                             Expr *E) {
@@ -9145,7 +9165,7 @@ static void ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D,
     handleSYCLKernelAttr(S, D, AL);
     break;
   case ParsedAttr::AT_SYCLSimd:
-    handleSimpleAttribute<SYCLSimdAttr>(S, D, AL);
+    handleSYCLSimdAttr(S, D, AL);
     break;
   case ParsedAttr::AT_SYCLDevice:
     handleSYCLDeviceAttr(S, D, AL);
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 5337ab3611ad3..700e0b05aa590 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -3521,10 +3521,14 @@ void Sema::MarkDevice(void) {
     }
   }
 
+
+  llvm::SmallPtrSet<FunctionDecl *, 10> AllKernelSets;
   for (Decl *D : syclDeviceDecls()) {
     if (auto SYCLKernel = dyn_cast<FunctionDecl>(D)) {
       llvm::SmallPtrSet<FunctionDecl *, 10> VisitedSet;
+      Marker.KernelSet.clear();
       Marker.CollectKernelSet(SYCLKernel, SYCLKernel, VisitedSet);
+      AllKernelSets.insert(Marker.KernelSet.begin(), Marker.KernelSet.end());
 
       // Let's propagate attributes from device functions to a SYCL kernels
       llvm::SmallVector<Attr *, 4> Attrs;
@@ -3534,7 +3538,7 @@ void Sema::MarkDevice(void) {
       FunctionDecl *KernelBody =
           Marker.CollectPossibleKernelAttributes(SYCLKernel, Attrs);
 
-      CheckSYCL2020Attributes(*this, SYCLKernel, VisitedSet);
+      CheckSYCL2020Attributes(*this, KernelBody, Marker.KernelSet);
 
       for (auto *A : Attrs) {
         switch (A->getKind()) {
@@ -3652,7 +3656,7 @@ void Sema::MarkDevice(void) {
       }
     }
   }
-  for (const auto &elt : Marker.KernelSet) {
+  for (const auto &elt : AllKernelSets) {
     if (FunctionDecl *Def = elt->getDefinition())
       Marker.TraverseStmt(Def->getBody());
   }
diff --git a/clang/test/SemaSYCL/sub-group-size.cpp b/clang/test/SemaSYCL/sub-group-size.cpp
index 0fad692581994..fa3c1068d4a1d 100644
--- a/clang/test/SemaSYCL/sub-group-size.cpp
+++ b/clang/test/SemaSYCL/sub-group-size.cpp
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -internal-isystem %S/Inputs -fsycl-is-device -internal-isystem %S/Inputs -fsyntax-only -verify %s
+// RUN: %clang_cc1 -internal-isystem %S/Inputs -fsycl-is-device -fsycl-explicit-simd -sycl-std=2020 -internal-isystem %S/Inputs -fsyntax-only -verify %s
 
 #include "Inputs/sycl.hpp"
 
@@ -42,3 +42,66 @@ void f6();
 // expected-warning@+1 {{'named_sub_group_size' attribute argument not supported: 'invalid'}}
 [[intel::named_sub_group_size(invalid)]]
 void f7();
+
+// expected-error@+2 {{'named_sub_group_size' and 'sycl_explicit_simd' attributes are not compatible}}
+// expected-note@+1 {{conflicting attribute is here}}
+[[intel::sycl_explicit_simd]][[intel::named_sub_group_size(automatic)]]
+void f8();
+// expected-error@+2 {{'sub_group_size' and 'sycl_explicit_simd' attributes are not compatible}}
+// expected-note@+1 {{conflicting attribute is here}}
+[[intel::sycl_explicit_simd]][[intel::sub_group_size(1)]]
+void f9();
+
+// expected-error@+1 {{'named_sub_group_size' and 'sycl_explicit_simd' attributes are not compatible}}
+[[intel::named_sub_group_size(primary)]]
+void f10();
+// expected-note@+1 {{conflicting attribute is here}}
+[[intel::sycl_explicit_simd]]
+void f10();
+
+void NoAttrFunc(){}
+SYCL_EXTERNAL void NoAttrExternalDefined() {}
+SYCL_EXTERNAL void NoAttrExternalNotDefined(); // #NoAttrExternalNotDefined
+
+struct Functor {
+  [[intel::named_sub_group_size(primary)]] void operator()() const {
+    NoAttrFunc();
+    NoAttrExternalDefined();
+    // expected-error@#NoAttrExternalNotDefined{{undefined 'SYCL_EXTERNAL' function must have a sub group size that matches the size specified for the kernel}}
+    // expected-note@-4{{conflicting attribute is here}}
+    NoAttrExternalNotDefined();
+  }
+};
+
+void calls_kernel_1() {
+  sycl::kernel_single_task<class Kernel1>([]() [[intel::named_sub_group_size(primary)]] {
+    NoAttrFunc();
+    NoAttrExternalDefined();
+    // expected-error@#NoAttrExternalNotDefined{{undefined 'SYCL_EXTERNAL' function must have a sub group size that matches the size specified for the kernel}}
+    // expected-note@-4{{conflicting attribute is here}}
+    NoAttrExternalNotDefined();
+  });
+}
+
+void calls_kernel_2() {
+  Functor F;
+  sycl::kernel_single_task<class Kernel2>(F);
+}
+
+// Func w/o attr called from kernel, kernel has attr.
+// normal func: fine
+// defined SYCL_EXTERNAL: fine
+// undef SYCL_EXTERNAL: Not fine
+// all are OK if kernel has 'default' attr.
+
+// Func w attr called from kernel, kernel has attr.
+// first matches default.
+// kernel matches default
+// both matches default
+// +SYCL_EXTERNAL
+
+// Func w attr called from kernel, kernel has no attr.
+// first matches default.
+// kernel matches default
+// both matches default
+// +SYCL_EXTERNAL

From 5260caa5c162203c8a170296aa334e9749700775 Mon Sep 17 00:00:00 2001
From: Erich Keane <erich.keane@intel.com>
Date: Tue, 30 Mar 2021 13:03:01 -0700
Subject: [PATCH 08/22] More progress on tests, found issue with defined
 SYCL_KERNEL functions

---
 clang/lib/Sema/SemaSYCL.cpp            | 12 +++++++-
 clang/test/SemaSYCL/sub-group-size.cpp | 39 ++++++++++++++++++++++----
 2 files changed, 44 insertions(+), 7 deletions(-)

diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 700e0b05aa590..90b0fb8d1df23 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -3476,8 +3476,16 @@ static void CheckSYCL2020SubGroupSizes(Sema &S, FunctionDecl *SYCLKernel,
 // on subsequent functions.
 static void
 CheckSYCL2020Attributes(Sema &S, FunctionDecl *SYCLKernel,
+                        FunctionDecl *KernelBody,
                         llvm::SmallPtrSetImpl<FunctionDecl *> &CalledFuncs) {
 
+  // SYCL_EXTERNAL functions don't have a body, but also include themselves in
+  // the call graph, so correct KernelBody and remove 'self' from the list.
+  if (!KernelBody) {
+    KernelBody = SYCLKernel;
+    CalledFuncs.erase(KernelBody);
+  }
+
   for (auto *FD : CalledFuncs) {
     for (auto *Attr : FD->attrs()) {
       switch (Attr->getKind()) {
@@ -3538,7 +3546,9 @@ void Sema::MarkDevice(void) {
       FunctionDecl *KernelBody =
           Marker.CollectPossibleKernelAttributes(SYCLKernel, Attrs);
 
-      CheckSYCL2020Attributes(*this, KernelBody, Marker.KernelSet);
+      // Check the list of implemented SYCL2020 attributes, which have different
+      // rules for propogation.
+      CheckSYCL2020Attributes(*this, SYCLKernel, KernelBody, Marker.KernelSet);
 
       for (auto *A : Attrs) {
         switch (A->getKind()) {
diff --git a/clang/test/SemaSYCL/sub-group-size.cpp b/clang/test/SemaSYCL/sub-group-size.cpp
index fa3c1068d4a1d..004f25e08df4b 100644
--- a/clang/test/SemaSYCL/sub-group-size.cpp
+++ b/clang/test/SemaSYCL/sub-group-size.cpp
@@ -1,4 +1,6 @@
-// RUN: %clang_cc1 -internal-isystem %S/Inputs -fsycl-is-device -fsycl-explicit-simd -sycl-std=2020 -internal-isystem %S/Inputs -fsyntax-only -verify %s
+// RUN: %clang_cc1 -internal-isystem %S/Inputs -fsycl-is-device -fsycl-explicit-simd -sycl-std=2020 -internal-isystem %S/Inputs -fsyntax-only -verify=expected,primary,integer %s
+// RUN: %clang_cc1 -internal-isystem %S/Inputs -fsycl-is-device -fsycl-explicit-simd -fsycl-default-sub-group-size=primary -sycl-std=2020 -internal-isystem %S/Inputs -fsyntax-only -verify=expected,integer %s
+// RUN: %clang_cc1 -internal-isystem %S/Inputs -fsycl-is-device -fsycl-explicit-simd -fsycl-default-sub-group-size=10 -sycl-std=2020 -internal-isystem %S/Inputs -fsyntax-only -verify=expected,primary %s
 
 #include "Inputs/sycl.hpp"
 
@@ -73,6 +75,8 @@ struct Functor {
   }
 };
 
+// If the kernel function has an attribute, only an undefined SYCL_EXTERNAL
+// should diagnose.
 void calls_kernel_1() {
   sycl::kernel_single_task<class Kernel1>([]() [[intel::named_sub_group_size(primary)]] {
     NoAttrFunc();
@@ -88,12 +92,35 @@ void calls_kernel_2() {
   sycl::kernel_single_task<class Kernel2>(F);
 }
 
-// Func w/o attr called from kernel, kernel has attr.
-// normal func: fine
-// defined SYCL_EXTERNAL: fine
-// undef SYCL_EXTERNAL: Not fine
-// all are OK if kernel has 'default' attr.
+// If the kernel doesn't have an attribute,
+[[intel::named_sub_group_size(primary)]] void AttrFunc(){} // #AttrFunc
+[[intel::named_sub_group_size(primary)]] SYCL_EXTERNAL void AttrExternalDefined() {} // #AttrExternalDefined
+[[intel::named_sub_group_size(primary)]] SYCL_EXTERNAL void AttrExternalNotDefined(); // #AttrExternalNotDefined
 
+void calls_kernel_3() {
+  sycl::kernel_single_task<class Kernel3>([]() {
+    // primary-error@#AttrFunc{{ :??}}
+    AttrFunc();
+    AttrExternalDefined();
+    // expected-error@#NoAttrExternalNotDefined{{undefined 'SYCL_EXTERNAL' function must have a sub group size that matches the size specified for the kernel}}
+    // expected-note@-4{{conflicting attribute is here}}
+    AttrExternalNotDefined();
+  });
+}
+
+[[intel::sub_group_size(4)]] void AttrFunc2(){}
+[[intel::sub_group_size(4)]] SYCL_EXTERNAL void AttrExternalDefined2() {}
+[[intel::sub_group_size(4)]] SYCL_EXTERNAL void AttrExternalNotDefined2();
+
+void calls_kernel_4() {
+  sycl::kernel_single_task<class Kernel4>([]() {
+    AttrFunc2();
+    AttrExternalDefined2();
+    // expected-error@#NoAttrExternalNotDefined{{undefined 'SYCL_EXTERNAL' function must have a sub group size that matches the size specified for the kernel}}
+    // expected-note@-4{{conflicting attribute is here}}
+    AttrExternalNotDefined2();
+  });
+}
 // Func w attr called from kernel, kernel has attr.
 // first matches default.
 // kernel matches default

From 10f2ec1312163a8db43f9d3fecce4549ce920ae0 Mon Sep 17 00:00:00 2001
From: Erich Keane <erich.keane@intel.com>
Date: Tue, 30 Mar 2021 13:53:39 -0700
Subject: [PATCH 09/22] Fix a few more semasycl diagnostics, fix command line
 arg

---
 clang/include/clang/Driver/Options.td     |  5 ++-
 clang/lib/Driver/ToolChains/Clang.cpp     |  4 ++
 clang/lib/Frontend/CompilerInvocation.cpp |  2 +-
 clang/lib/Sema/SemaSYCL.cpp               | 14 ++++---
 clang/test/SemaSYCL/sub-group-size.cpp    | 51 +++++++++++++----------
 5 files changed, 46 insertions(+), 30 deletions(-)

diff --git a/clang/include/clang/Driver/Options.td b/clang/include/clang/Driver/Options.td
index e1cab4547c484..da73d12a606db 100644
--- a/clang/include/clang/Driver/Options.td
+++ b/clang/include/clang/Driver/Options.td
@@ -5548,10 +5548,11 @@ def sycl_std_EQ : Joined<["-"], "sycl-std=">, Group<sycl_Group>,
 
 def fsycl_default_sub_group_size
     : Separate<["-"], "fsycl-default-sub-group-size">,
-      HelpText<"Set the default sub group size for SYCL kernels">;
+      HelpText<"Set the default sub group size for SYCL kernels">,
+      Flags<[CC1Option]>;
 def fsycl_default_sub_group_size_EQ
     : Joined<["-"], "fsycl-default-sub-group-size=">,
-      Alias<fsycl_default_sub_group_size>;
+      Alias<fsycl_default_sub_group_size>, Flags<[CC1Option]>;
 
 defm cuda_approx_transcendentals : BoolFOption<"cuda-approx-transcendentals",
   LangOpts<"CUDADeviceApproxTranscendentals">, DefaultFalse,
diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp
index 7ecd402371704..97b028da37f33 100644
--- a/clang/lib/Driver/ToolChains/Clang.cpp
+++ b/clang/lib/Driver/ToolChains/Clang.cpp
@@ -5881,6 +5881,10 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA,
   if (Args.hasFlag(options::OPT_fsycl, options::OPT_fno_sycl, false))
     Args.AddLastArg(CmdArgs, options::OPT_sycl_std_EQ);
 
+  // Forward -fsycl-default-sub-group-size if in SYCL mode.
+  if (Args.hasFlag(options::OPT_fsycl, options::OPT_fno_sycl, false))
+    Args.AddLastArg(CmdArgs, options::OPT_fsycl_default_sub_group_size);
+
   // Forward -fsycl-instrument-device-code option to cc1. This option can only
   // be used with spir triple.
   if (Arg *A = Args.getLastArg(options::OPT_fsycl_instrument_device_code)) {
diff --git a/clang/lib/Frontend/CompilerInvocation.cpp b/clang/lib/Frontend/CompilerInvocation.cpp
index b79b06686ca2b..90aa6e8c32b67 100644
--- a/clang/lib/Frontend/CompilerInvocation.cpp
+++ b/clang/lib/Frontend/CompilerInvocation.cpp
@@ -3740,7 +3740,7 @@ bool CompilerInvocation::ParseLangArgsImpl(LangOptions &Opts, ArgList &Args,
     if (Opts.getDefaultSubGroupSizeType() ==
         LangOptions::SubGroupSizeType::Integer) {
       int64_t IntResult;
-      if (Value.getAsInteger(10, IntResult)) {
+      if (!Value.getAsInteger(10, IntResult)) {
         Opts.DefaultSubGroupSize = IntResult;
       } else {
         Diags.Report(diag::err_drv_invalid_value)
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 90b0fb8d1df23..bb81eafa3107b 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -3478,14 +3478,16 @@ static void
 CheckSYCL2020Attributes(Sema &S, FunctionDecl *SYCLKernel,
                         FunctionDecl *KernelBody,
                         llvm::SmallPtrSetImpl<FunctionDecl *> &CalledFuncs) {
+  // We don't care about the kernel itself or the body, so remove them from the
+  // list.
+  CalledFuncs.erase(SYCLKernel);
+  CalledFuncs.erase(KernelBody);
 
-  // SYCL_EXTERNAL functions don't have a body, but also include themselves in
-  // the call graph, so correct KernelBody and remove 'self' from the list.
-  if (!KernelBody) {
-    KernelBody = SYCLKernel;
-    CalledFuncs.erase(KernelBody);
+  // If the kernel has a body, we should get the attributes for the kernel from
+  // there instead, so that we get the functor object.
+  if (KernelBody) {
+    SYCLKernel = KernelBody;
   }
-
   for (auto *FD : CalledFuncs) {
     for (auto *Attr : FD->attrs()) {
       switch (Attr->getKind()) {
diff --git a/clang/test/SemaSYCL/sub-group-size.cpp b/clang/test/SemaSYCL/sub-group-size.cpp
index 004f25e08df4b..145cb79e87a6e 100644
--- a/clang/test/SemaSYCL/sub-group-size.cpp
+++ b/clang/test/SemaSYCL/sub-group-size.cpp
@@ -65,20 +65,10 @@ void NoAttrFunc(){}
 SYCL_EXTERNAL void NoAttrExternalDefined() {}
 SYCL_EXTERNAL void NoAttrExternalNotDefined(); // #NoAttrExternalNotDefined
 
-struct Functor {
-  [[intel::named_sub_group_size(primary)]] void operator()() const {
-    NoAttrFunc();
-    NoAttrExternalDefined();
-    // expected-error@#NoAttrExternalNotDefined{{undefined 'SYCL_EXTERNAL' function must have a sub group size that matches the size specified for the kernel}}
-    // expected-note@-4{{conflicting attribute is here}}
-    NoAttrExternalNotDefined();
-  }
-};
-
 // If the kernel function has an attribute, only an undefined SYCL_EXTERNAL
 // should diagnose.
 void calls_kernel_1() {
-  sycl::kernel_single_task<class Kernel1>([]() [[intel::named_sub_group_size(primary)]] {
+  sycl::kernel_single_task<class Kernel1>([]() [[intel::named_sub_group_size(automatic)]] {
     NoAttrFunc();
     NoAttrExternalDefined();
     // expected-error@#NoAttrExternalNotDefined{{undefined 'SYCL_EXTERNAL' function must have a sub group size that matches the size specified for the kernel}}
@@ -87,6 +77,16 @@ void calls_kernel_1() {
   });
 }
 
+struct Functor {
+  [[intel::named_sub_group_size(automatic)]] void operator()() const {
+    NoAttrFunc();
+ //   NoAttrExternalDefined();
+    // expected-error@#NoAttrExternalNotDefined{{undefined 'SYCL_EXTERNAL' function must have a sub group size that matches the size specified for the kernel}}
+    // expected-note@-4{{conflicting attribute is here}}
+    NoAttrExternalNotDefined();
+  }
+};
+
 void calls_kernel_2() {
   Functor F;
   sycl::kernel_single_task<class Kernel2>(F);
@@ -98,29 +98,38 @@ void calls_kernel_2() {
 [[intel::named_sub_group_size(primary)]] SYCL_EXTERNAL void AttrExternalNotDefined(); // #AttrExternalNotDefined
 
 void calls_kernel_3() {
-  sycl::kernel_single_task<class Kernel3>([]() {
-    // primary-error@#AttrFunc{{ :??}}
+  sycl::kernel_single_task<class Kernel3>([]() { // #Kernel3
+    // primary-error@#AttrFunc{{kernel-called function must have a sub group size that matches the size specified for the kernel}}
+    // primary-note@#Kernel3{{kernel declared here}}
     AttrFunc();
+    // primary-error@#AttrExternalDefined{{kernel-called function must have a sub group size that matches the size specified for the kernel}}
+    // primary-note@#Kernel3{{kernel declared here}}
     AttrExternalDefined();
-    // expected-error@#NoAttrExternalNotDefined{{undefined 'SYCL_EXTERNAL' function must have a sub group size that matches the size specified for the kernel}}
-    // expected-note@-4{{conflicting attribute is here}}
+    // primary-error@#AttrExternalNotDefined{{kernel-called function must have a sub group size that matches the size specified for the kernel}}
+    // primary-note@#Kernel3{{kernel declared here}}
     AttrExternalNotDefined();
   });
 }
 
-[[intel::sub_group_size(4)]] void AttrFunc2(){}
-[[intel::sub_group_size(4)]] SYCL_EXTERNAL void AttrExternalDefined2() {}
-[[intel::sub_group_size(4)]] SYCL_EXTERNAL void AttrExternalNotDefined2();
+[[intel::sub_group_size(10)]] void AttrFunc2(){} // #AttrFunc2
+[[intel::sub_group_size(10)]] SYCL_EXTERNAL void AttrExternalDefined2() {} // #AttrExternalDefined2
+[[intel::sub_group_size(10)]] SYCL_EXTERNAL void AttrExternalNotDefined2(); // #AttrExternalNotDefined2
 
 void calls_kernel_4() {
-  sycl::kernel_single_task<class Kernel4>([]() {
+  sycl::kernel_single_task<class Kernel4>([]() { // #Kernel4
+    // integer-error@#AttrFunc2{{kernel-called function must have a sub group size that matches the size specified for the kernel}}
+    // integer-note@#Kernel4{{kernel declared here}}
     AttrFunc2();
+    // integer-error@#AttrExternalDefined2{{kernel-called function must have a sub group size that matches the size specified for the kernel}}
+    // integer-note@#Kernel4{{kernel declared here}}
     AttrExternalDefined2();
-    // expected-error@#NoAttrExternalNotDefined{{undefined 'SYCL_EXTERNAL' function must have a sub group size that matches the size specified for the kernel}}
-    // expected-note@-4{{conflicting attribute is here}}
+    // integer-error@#AttrExternalNotDefined2{{kernel-called function must have a sub group size that matches the size specified for the kernel}}
+    // integer-note@#Kernel4{{kernel declared here}}
     AttrExternalNotDefined2();
   });
 }
+
+// TODO!
 // Func w attr called from kernel, kernel has attr.
 // first matches default.
 // kernel matches default

From 22671a13c38fcf92baa4270806f34bffbd057329 Mon Sep 17 00:00:00 2001
From: Erich Keane <erich.keane@intel.com>
Date: Wed, 31 Mar 2021 05:55:43 -0700
Subject: [PATCH 10/22] Get to a reasonable position so that we can pick this
 up later.

Currently the discussion on the design has lead me to believe that
writing this patch is premature. Some of the in-flight changes to the
spec will severely change what this implementation looks like, so we're
going to put this patch on hold until the spec settles.

This is hopefully in a reasonable enough position where we can pick it
up later without trouble.
---
 clang/lib/CodeGen/CodeGenFunction.cpp  | 2 ++
 clang/lib/Sema/SemaSYCL.cpp            | 4 ++++
 clang/test/SemaSYCL/sub-group-size.cpp | 7 ++++++-
 3 files changed, 12 insertions(+), 1 deletion(-)

diff --git a/clang/lib/CodeGen/CodeGenFunction.cpp b/clang/lib/CodeGen/CodeGenFunction.cpp
index bcbeaf9da7e08..bd4470e01040d 100644
--- a/clang/lib/CodeGen/CodeGenFunction.cpp
+++ b/clang/lib/CodeGen/CodeGenFunction.cpp
@@ -645,6 +645,8 @@ void CodeGenFunction::EmitOpenCLKernelMetadata(const FunctionDecl *FD,
   const IntelReqdSubGroupSizeAttr *ReqSubGroup =
       FD->getAttr<IntelReqdSubGroupSizeAttr>();
 
+  // TODO: This section and the next (sub-group-size) is in flux and awaiting
+  // a review of the spec.
   // To support the SYCL 2020 spelling with no propagation, only emit for
   // kernel-or-device when that spelling, fall-back to old behavior.
   if (ReqSubGroup && (IsKernelOrDevice || !ReqSubGroup->isSYCL2020Spelling())) {
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index bb81eafa3107b..e9e9190e66cf4 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -3487,7 +3487,11 @@ CheckSYCL2020Attributes(Sema &S, FunctionDecl *SYCLKernel,
   // there instead, so that we get the functor object.
   if (KernelBody) {
     SYCLKernel = KernelBody;
+    // TODO: In the else case here, we likely want to propagate kernel-body
+    // attributes to the kernel, so that we can limit codegen of these to just
+    //  the kernel.
   }
+
   for (auto *FD : CalledFuncs) {
     for (auto *Attr : FD->attrs()) {
       switch (Attr->getKind()) {
diff --git a/clang/test/SemaSYCL/sub-group-size.cpp b/clang/test/SemaSYCL/sub-group-size.cpp
index 145cb79e87a6e..dca7126a98eff 100644
--- a/clang/test/SemaSYCL/sub-group-size.cpp
+++ b/clang/test/SemaSYCL/sub-group-size.cpp
@@ -129,15 +129,20 @@ void calls_kernel_4() {
   });
 }
 
-// TODO!
+// TODO, test the following:
 // Func w attr called from kernel, kernel has attr.
+// Neither matches default.
 // first matches default.
 // kernel matches default
 // both matches default
 // +SYCL_EXTERNAL
 
 // Func w attr called from kernel, kernel has no attr.
+// Neither matches default.
 // first matches default.
 // kernel matches default
 // both matches default
 // +SYCL_EXTERNAL
+
+
+// TODO: CodeGen tests

From 5036d3a686ecd9d578927b2785e69c28bb4f4a2b Mon Sep 17 00:00:00 2001
From: Erich Keane <erich.keane@intel.com>
Date: Fri, 9 Apr 2021 07:52:41 -0700
Subject: [PATCH 11/22] Allow string parameters, finish up the sema tests.

Still need to do the code gen, but that is broken at least until the
rebase on #3475, which is awaiting merge.
---
 clang/lib/CodeGen/CodeGenFunction.cpp     |  2 -
 clang/lib/Sema/SemaDeclAttr.cpp           | 19 ++++----
 clang/test/CodeGenSYCL/sub-group-size.cpp | 11 +++++
 clang/test/SemaSYCL/sub-group-size.cpp    | 56 +++++++++++++++--------
 4 files changed, 57 insertions(+), 31 deletions(-)
 create mode 100644 clang/test/CodeGenSYCL/sub-group-size.cpp

diff --git a/clang/lib/CodeGen/CodeGenFunction.cpp b/clang/lib/CodeGen/CodeGenFunction.cpp
index eab6acda6a83c..1c09702a07504 100644
--- a/clang/lib/CodeGen/CodeGenFunction.cpp
+++ b/clang/lib/CodeGen/CodeGenFunction.cpp
@@ -653,8 +653,6 @@ void CodeGenFunction::EmitOpenCLKernelMetadata(const FunctionDecl *FD,
   const IntelReqdSubGroupSizeAttr *ReqSubGroup =
       FD->getAttr<IntelReqdSubGroupSizeAttr>();
 
-  // TODO: This section and the next (sub-group-size) is in flux and awaiting
-  // a review of the spec.
   // To support the SYCL 2020 spelling with no propagation, only emit for
   // kernel-or-device when that spelling, fall-back to old behavior.
   if (ReqSubGroup && (IsKernelOrDevice || !ReqSubGroup->isSYCL2020Spelling())) {
diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index f8ae2c42109c6..ec965078c39b9 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -3245,20 +3245,21 @@ static void handleIntelNamedSubGroupSize(Sema &S, Decl *D,
   if (checkAttrMutualExclusion<SYCLSimdAttr>(S, D, AL))
     return;
 
-  if (!AL.isArgIdent(0)) {
-    S.Diag(AL.getArgAsExpr(0)->getBeginLoc(), diag::err_attribute_argument_type)
-        << AL << AANT_ArgumentIdentifier;
+  StringRef SizeStr;
+  SourceLocation Loc;
+  if (AL.isArgIdent(0)) {
+    IdentifierLoc *IL = AL.getArgAsIdent(0);
+    SizeStr = IL->Ident->getName();
+    Loc = IL->Loc;
+  } else if (!S.checkStringLiteralArgumentAttr(AL, 0, SizeStr, &Loc)) {
     return;
   }
 
   IntelNamedSubGroupSizeAttr::SubGroupSizeType SizeType;
-  IdentifierLoc *IL = AL.getArgAsIdent(0);
-  if (!IntelNamedSubGroupSizeAttr::ConvertStrToSubGroupSizeType(
-          IL->Ident->getName(), SizeType)) {
-    S.Diag(IL->Loc, diag::warn_attribute_type_not_supported) << AL << IL->Ident;
-    return;
+  if (!IntelNamedSubGroupSizeAttr::ConvertStrToSubGroupSizeType(SizeStr,
+                                                                SizeType)) {
+    S.Diag(Loc, diag::warn_attribute_type_not_supported) << AL << SizeStr;
   }
-
   D->addAttr(IntelNamedSubGroupSizeAttr::Create(S.Context, SizeType, AL));
 }
 
diff --git a/clang/test/CodeGenSYCL/sub-group-size.cpp b/clang/test/CodeGenSYCL/sub-group-size.cpp
new file mode 100644
index 0000000000000..911566fb4a439
--- /dev/null
+++ b/clang/test/CodeGenSYCL/sub-group-size.cpp
@@ -0,0 +1,11 @@
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2020 -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s --check-prefix=NONE
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -fsycl-default-sub-group-size=primary -sycl-std=2020 -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s --check-prefix=PRIM_DEF
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -fsycl-default-sub-group-size=10 -sycl-std=2020 -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s --check-prefix=TEN_DEF
+
+#include "Inputs/sycl.hpp"
+using namespace cl::sycl;
+
+void default_behavior() {
+  kernel_single_task<class Kernel1>([]() {
+  });
+}
diff --git a/clang/test/SemaSYCL/sub-group-size.cpp b/clang/test/SemaSYCL/sub-group-size.cpp
index dca7126a98eff..d2b86d7220365 100644
--- a/clang/test/SemaSYCL/sub-group-size.cpp
+++ b/clang/test/SemaSYCL/sub-group-size.cpp
@@ -1,6 +1,6 @@
-// RUN: %clang_cc1 -internal-isystem %S/Inputs -fsycl-is-device -fsycl-explicit-simd -sycl-std=2020 -internal-isystem %S/Inputs -fsyntax-only -verify=expected,primary,integer %s
-// RUN: %clang_cc1 -internal-isystem %S/Inputs -fsycl-is-device -fsycl-explicit-simd -fsycl-default-sub-group-size=primary -sycl-std=2020 -internal-isystem %S/Inputs -fsyntax-only -verify=expected,integer %s
-// RUN: %clang_cc1 -internal-isystem %S/Inputs -fsycl-is-device -fsycl-explicit-simd -fsycl-default-sub-group-size=10 -sycl-std=2020 -internal-isystem %S/Inputs -fsyntax-only -verify=expected,primary %s
+// RUN: %clang_cc1 -internal-isystem %S/Inputs -fsycl-is-device -sycl-std=2020 -internal-isystem %S/Inputs -fsyntax-only -verify=expected,primary,integer %s
+// RUN: %clang_cc1 -internal-isystem %S/Inputs -fsycl-is-device -fsycl-default-sub-group-size=primary -sycl-std=2020 -internal-isystem %S/Inputs -fsyntax-only -verify=expected,integer %s
+// RUN: %clang_cc1 -internal-isystem %S/Inputs -fsycl-is-device -fsycl-default-sub-group-size=10 -sycl-std=2020 -internal-isystem %S/Inputs -fsyntax-only -verify=expected,primary %s
 
 #include "Inputs/sycl.hpp"
 
@@ -41,7 +41,7 @@ void f6();
 [[intel::named_sub_group_size(automatic)]]
 void f6();
 
-// expected-warning@+1 {{'named_sub_group_size' attribute argument not supported: 'invalid'}}
+// expected-warning@+1 {{'named_sub_group_size' attribute argument not supported: invalid}}
 [[intel::named_sub_group_size(invalid)]]
 void f7();
 
@@ -61,6 +61,24 @@ void f10();
 [[intel::sycl_explicit_simd]]
 void f10();
 
+// expected-error@+1 {{'named_sub_group_size' and 'sycl_explicit_simd' attributes are not compatible}}
+[[intel::named_sub_group_size("primary")]]
+void f11();
+// expected-note@+1 {{conflicting attribute is here}}
+[[intel::sycl_explicit_simd]]
+void f11();
+
+// expected-error@+1 {{'named_sub_group_size' and 'sycl_explicit_simd' attributes are not compatible}}
+[[intel::named_sub_group_size("automatic")]]
+void f12();
+// expected-note@+1 {{conflicting attribute is here}}
+[[intel::sycl_explicit_simd]]
+void f12();
+
+// expected-warning@+1 {{'named_sub_group_size' attribute argument not supported: invalid string}}
+[[intel::named_sub_group_size("invalid string")]]
+void f13();
+
 void NoAttrFunc(){}
 SYCL_EXTERNAL void NoAttrExternalDefined() {}
 SYCL_EXTERNAL void NoAttrExternalNotDefined(); // #NoAttrExternalNotDefined
@@ -129,20 +147,18 @@ void calls_kernel_4() {
   });
 }
 
-// TODO, test the following:
-// Func w attr called from kernel, kernel has attr.
-// Neither matches default.
-// first matches default.
-// kernel matches default
-// both matches default
-// +SYCL_EXTERNAL
-
-// Func w attr called from kernel, kernel has no attr.
-// Neither matches default.
-// first matches default.
-// kernel matches default
-// both matches default
-// +SYCL_EXTERNAL
-
+// Both have an attribute.
+void calls_kernel_5() {
+  sycl::kernel_single_task<class Kernel5>([]() [[intel::named_sub_group_size(automatic)]] { // #Kernel5
+    // expected-error@#AttrFunc{{kernel-called function must have a sub group size that matches the size specified for the kernel}}
+    // expected-note@#Kernel5{{conflicting attribute is here}}
+    AttrFunc();
+    // expected-error@#AttrExternalDefined{{kernel-called function must have a sub group size that matches the size specified for the kernel}}
+    // expected-note@#Kernel5{{conflicting attribute is here}}
+    AttrExternalDefined();
+    // expected-error@#AttrExternalNotDefined{{kernel-called function must have a sub group size that matches the size specified for the kernel}}
+    // expected-note@#Kernel5{{conflicting attribute is here}}
+    AttrExternalNotDefined();
 
-// TODO: CodeGen tests
+  });
+}

From 0318b402cde20a93e604d78132f04a8eddc60c7c Mon Sep 17 00:00:00 2001
From: Erich Keane <erich.keane@intel.com>
Date: Fri, 9 Apr 2021 09:31:59 -0700
Subject: [PATCH 12/22] Clang-format fixes

---
 clang/test/SemaSYCL/sub-group-size.cpp | 15 +++++----------
 1 file changed, 5 insertions(+), 10 deletions(-)

diff --git a/clang/test/SemaSYCL/sub-group-size.cpp b/clang/test/SemaSYCL/sub-group-size.cpp
index d2b86d7220365..ffaa6ee5dbb2c 100644
--- a/clang/test/SemaSYCL/sub-group-size.cpp
+++ b/clang/test/SemaSYCL/sub-group-size.cpp
@@ -62,22 +62,17 @@ void f10();
 void f10();
 
 // expected-error@+1 {{'named_sub_group_size' and 'sycl_explicit_simd' attributes are not compatible}}
-[[intel::named_sub_group_size("primary")]]
-void f11();
+[[intel::named_sub_group_size("primary")]] void f11();
 // expected-note@+1 {{conflicting attribute is here}}
-[[intel::sycl_explicit_simd]]
-void f11();
+[[intel::sycl_explicit_simd]] void f11();
 
 // expected-error@+1 {{'named_sub_group_size' and 'sycl_explicit_simd' attributes are not compatible}}
-[[intel::named_sub_group_size("automatic")]]
-void f12();
+[[intel::named_sub_group_size("automatic")]] void f12();
 // expected-note@+1 {{conflicting attribute is here}}
-[[intel::sycl_explicit_simd]]
-void f12();
+[[intel::sycl_explicit_simd]] void f12();
 
 // expected-warning@+1 {{'named_sub_group_size' attribute argument not supported: invalid string}}
-[[intel::named_sub_group_size("invalid string")]]
-void f13();
+[[intel::named_sub_group_size("invalid string")]] void f13();
 
 void NoAttrFunc(){}
 SYCL_EXTERNAL void NoAttrExternalDefined() {}

From 612ea7863e96cb2c4735d1a7b30921bc513e8979 Mon Sep 17 00:00:00 2001
From: Erich Keane <erich.keane@intel.com>
Date: Fri, 9 Apr 2021 09:37:09 -0700
Subject: [PATCH 13/22] more clang-format fixes

---
 clang/lib/Sema/SemaDeclAttr.cpp        |  4 +-
 clang/lib/Sema/SemaSYCL.cpp            |  1 -
 clang/test/SemaSYCL/sub-group-size.cpp | 57 ++++++++++----------------
 3 files changed, 22 insertions(+), 40 deletions(-)

diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index ec965078c39b9..5c9df66894892 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -3216,7 +3216,6 @@ static void handleIntelReqdSubGroupSize(Sema &S, Decl *D,
   S.AddIntelReqdSubGroupSize(D, AL, E);
 }
 
-
 IntelNamedSubGroupSizeAttr *
 Sema::MergeIntelNamedSubGroupSizeAttr(Decl *D,
                                       const IntelNamedSubGroupSizeAttr &A) {
@@ -3272,8 +3271,7 @@ SYCLSimdAttr *Sema::MergeSYCLSimdAttr(Decl *D, const SYCLSimdAttr &A) {
   return A.clone(Context);
 }
 
-static void handleSYCLSimdAttr(Sema &S, Decl *D,
-                                         const ParsedAttr &AL) {
+static void handleSYCLSimdAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
   if (checkAttrMutualExclusion<IntelReqdSubGroupSizeAttr>(S, D, AL))
     return;
   if (checkAttrMutualExclusion<IntelNamedSubGroupSizeAttr>(S, D, AL))
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 567c280153396..ff119aed62758 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -3619,7 +3619,6 @@ void Sema::MarkDevice(void) {
     }
   }
 
-
   llvm::SmallPtrSet<FunctionDecl *, 10> AllKernelSets;
   for (Decl *D : syclDeviceDecls()) {
     if (auto SYCLKernel = dyn_cast<FunctionDecl>(D)) {
diff --git a/clang/test/SemaSYCL/sub-group-size.cpp b/clang/test/SemaSYCL/sub-group-size.cpp
index ffaa6ee5dbb2c..2a955def64b5f 100644
--- a/clang/test/SemaSYCL/sub-group-size.cpp
+++ b/clang/test/SemaSYCL/sub-group-size.cpp
@@ -6,60 +6,45 @@
 
 // expected-error@+2 {{'named_sub_group_size' and 'sub_group_size' attributes are not compatible}}
 // expected-note@+1 {{conflicting attribute is here}}
-[[intel::sub_group_size(1)]][[intel::named_sub_group_size(automatic)]]
-void f1();
+[[intel::sub_group_size(1)]] [[intel::named_sub_group_size(automatic)]] void f1();
 // expected-error@+2 {{'sub_group_size' and 'named_sub_group_size' attributes are not compatible}}
 // expected-note@+1 {{conflicting attribute is here}}
-[[intel::named_sub_group_size(primary)]][[intel::sub_group_size(1)]]
-void f2();
+[[intel::named_sub_group_size(primary)]] [[intel::sub_group_size(1)]] void f2();
 
 // expected-error@+1 {{'sub_group_size' and 'named_sub_group_size' attributes are not compatible}}
-[[intel::sub_group_size(1)]]
-void f3();
+[[intel::sub_group_size(1)]] void f3();
 // expected-note@+1 {{conflicting attribute is here}}
-[[intel::named_sub_group_size(primary)]]
-void f3();
+[[intel::named_sub_group_size(primary)]] void f3();
 
 // expected-error@+1 {{'named_sub_group_size' and 'sub_group_size' attributes are not compatible}}
-[[intel::named_sub_group_size(primary)]]
-void f4();
+[[intel::named_sub_group_size(primary)]] void f4();
 // expected-note@+1 {{conflicting attribute is here}}
-[[intel::sub_group_size(1)]]
-void f4();
+[[intel::sub_group_size(1)]] void f4();
 
 // expected-note@+1 {{previous attribute is here}}
-[[intel::named_sub_group_size(automatic)]]
-void f5();
+[[intel::named_sub_group_size(automatic)]] void f5();
 
 // expected-warning@+1 {{attribute 'named_sub_group_size' is already applied with different arguments}}
-[[intel::named_sub_group_size(primary)]]
-void f5();
+[[intel::named_sub_group_size(primary)]] void f5();
 
-[[intel::named_sub_group_size(automatic)]]
-void f6();
+[[intel::named_sub_group_size(automatic)]] void f6();
 
-[[intel::named_sub_group_size(automatic)]]
-void f6();
+[[intel::named_sub_group_size(automatic)]] void f6();
 
 // expected-warning@+1 {{'named_sub_group_size' attribute argument not supported: invalid}}
-[[intel::named_sub_group_size(invalid)]]
-void f7();
+[[intel::named_sub_group_size(invalid)]] void f7();
 
 // expected-error@+2 {{'named_sub_group_size' and 'sycl_explicit_simd' attributes are not compatible}}
 // expected-note@+1 {{conflicting attribute is here}}
-[[intel::sycl_explicit_simd]][[intel::named_sub_group_size(automatic)]]
-void f8();
+[[intel::sycl_explicit_simd]] [[intel::named_sub_group_size(automatic)]] void f8();
 // expected-error@+2 {{'sub_group_size' and 'sycl_explicit_simd' attributes are not compatible}}
 // expected-note@+1 {{conflicting attribute is here}}
-[[intel::sycl_explicit_simd]][[intel::sub_group_size(1)]]
-void f9();
+[[intel::sycl_explicit_simd]] [[intel::sub_group_size(1)]] void f9();
 
 // expected-error@+1 {{'named_sub_group_size' and 'sycl_explicit_simd' attributes are not compatible}}
-[[intel::named_sub_group_size(primary)]]
-void f10();
+[[intel::named_sub_group_size(primary)]] void f10();
 // expected-note@+1 {{conflicting attribute is here}}
-[[intel::sycl_explicit_simd]]
-void f10();
+[[intel::sycl_explicit_simd]] void f10();
 
 // expected-error@+1 {{'named_sub_group_size' and 'sycl_explicit_simd' attributes are not compatible}}
 [[intel::named_sub_group_size("primary")]] void f11();
@@ -74,7 +59,7 @@ void f10();
 // expected-warning@+1 {{'named_sub_group_size' attribute argument not supported: invalid string}}
 [[intel::named_sub_group_size("invalid string")]] void f13();
 
-void NoAttrFunc(){}
+void NoAttrFunc() {}
 SYCL_EXTERNAL void NoAttrExternalDefined() {}
 SYCL_EXTERNAL void NoAttrExternalNotDefined(); // #NoAttrExternalNotDefined
 
@@ -93,7 +78,7 @@ void calls_kernel_1() {
 struct Functor {
   [[intel::named_sub_group_size(automatic)]] void operator()() const {
     NoAttrFunc();
- //   NoAttrExternalDefined();
+    //   NoAttrExternalDefined();
     // expected-error@#NoAttrExternalNotDefined{{undefined 'SYCL_EXTERNAL' function must have a sub group size that matches the size specified for the kernel}}
     // expected-note@-4{{conflicting attribute is here}}
     NoAttrExternalNotDefined();
@@ -106,8 +91,8 @@ void calls_kernel_2() {
 }
 
 // If the kernel doesn't have an attribute,
-[[intel::named_sub_group_size(primary)]] void AttrFunc(){} // #AttrFunc
-[[intel::named_sub_group_size(primary)]] SYCL_EXTERNAL void AttrExternalDefined() {} // #AttrExternalDefined
+[[intel::named_sub_group_size(primary)]] void AttrFunc() {}                           // #AttrFunc
+[[intel::named_sub_group_size(primary)]] SYCL_EXTERNAL void AttrExternalDefined() {}  // #AttrExternalDefined
 [[intel::named_sub_group_size(primary)]] SYCL_EXTERNAL void AttrExternalNotDefined(); // #AttrExternalNotDefined
 
 void calls_kernel_3() {
@@ -124,8 +109,8 @@ void calls_kernel_3() {
   });
 }
 
-[[intel::sub_group_size(10)]] void AttrFunc2(){} // #AttrFunc2
-[[intel::sub_group_size(10)]] SYCL_EXTERNAL void AttrExternalDefined2() {} // #AttrExternalDefined2
+[[intel::sub_group_size(10)]] void AttrFunc2() {}                           // #AttrFunc2
+[[intel::sub_group_size(10)]] SYCL_EXTERNAL void AttrExternalDefined2() {}  // #AttrExternalDefined2
 [[intel::sub_group_size(10)]] SYCL_EXTERNAL void AttrExternalNotDefined2(); // #AttrExternalNotDefined2
 
 void calls_kernel_4() {

From b692be120c68ce18726f759f5f82a8e18f3d4c78 Mon Sep 17 00:00:00 2001
From: Erich Keane <erich.keane@intel.com>
Date: Mon, 12 Apr 2021 08:33:49 -0700
Subject: [PATCH 14/22] Aaron's review comments & Merge build failure fixes

---
 clang/include/clang/Basic/Attr.td         |  6 ++-
 clang/include/clang/Basic/AttrDocs.td     | 50 ++++++++++++-----------
 clang/include/clang/Basic/LangOptions.def |  2 +-
 clang/include/clang/Sema/Sema.h           |  1 -
 clang/lib/Sema/SemaDecl.cpp               |  2 -
 clang/lib/Sema/SemaDeclAttr.cpp           | 40 +-----------------
 clang/lib/Sema/SemaSYCL.cpp               |  8 +++-
 clang/test/SemaSYCL/sub-group-size.cpp    | 20 ++++-----
 8 files changed, 49 insertions(+), 80 deletions(-)

diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td
index 0728ec30d4bd1..04fd3e64ea744 100644
--- a/clang/include/clang/Basic/Attr.td
+++ b/clang/include/clang/Basic/Attr.td
@@ -1213,6 +1213,7 @@ def SYCLSimd : InheritableAttr {
   let Subjects = SubjectList<[Function, GlobalVar]>;
   let Documentation = [SYCLSimdDocs];
   let SupportsNonconformingLambdaSyntax = 1;
+  let SimpleHandler = 1;
 }
 
 // Available in SYCL explicit SIMD extension. Binds a file scope private
@@ -1441,7 +1442,7 @@ def IntelReqdSubGroupSize: InheritableAttr {
   let LangOpts = [OpenCL, SYCLIsDevice, SilentlyIgnoreSYCLIsHost];
   let SupportsNonconformingLambdaSyntax = 1;
   let Accessors =
-      [Accessor<"isSYCL2020Spelling", [CXX11<"intel", "sub_group_size">]>];
+    [Accessor<"isSYCL2020Spelling", [CXX11<"intel", "sub_group_size">]>];
 }
 
 def IntelNamedSubGroupSize : InheritableAttr {
@@ -1454,6 +1455,9 @@ def IntelNamedSubGroupSize : InheritableAttr {
   let SupportsNonconformingLambdaSyntax = 1;
 }
 
+def : MutualExclusions<
+          [IntelReqdSubGroupSize, IntelNamedSubGroupSize, SYCLSimd]>;
+
 // This attribute is both a type attribute, and a declaration attribute (for
 // parameter variables).
 def OpenCLAccess : Attr {
diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td
index 1420c53b3d061..5978c7a4573f2 100644
--- a/clang/include/clang/Basic/AttrDocs.td
+++ b/clang/include/clang/Basic/AttrDocs.td
@@ -453,16 +453,16 @@ def NoMergeDocs : Documentation {
   let Category = DocCatStmt;
   let Content = [{
 If a statement is marked ``nomerge`` and contains call expressions, those call
-expressions inside the statement will not be merged during optimization. This 
+expressions inside the statement will not be merged during optimization. This
 attribute can be used to prevent the optimizer from obscuring the source
 location of certain calls. For example, it will prevent tail merging otherwise
 identical code sequences that raise an exception or terminate the program. Tail
 merging normally reduces the precision of source location information, making
 stack traces less useful for debugging. This attribute gives the user control
-over the tradeoff between code size and debug information precision. 
+over the tradeoff between code size and debug information precision.
 
-``nomerge`` attribute can also be used as function attribute to prevent all 
-calls to the specified function from merging. It has no effect on indirect 
+``nomerge`` attribute can also be used as function attribute to prevent all
+calls to the specified function from merging. It has no effect on indirect
 calls.
   }];
 }
@@ -2544,11 +2544,11 @@ argument to **clEnqueueNDRangeKernel** (in OpenCL) or to
 generated code appropriately for the kernel to which attribute is applied.
 
 While semantic of this attribute is the same between OpenCL and SYCL,
-spelling is a bit different: 
+spelling is a bit different:
 
 SYCL 1.2.1 describes ``[[cl::reqd_work_group_size(X, Y, Z)]]`` spelling: this
-attribute is legal on device functions and is propagated down to any caller of 
-those device functions, such that the kernel attributes are the sum of all 
+attribute is legal on device functions and is propagated down to any caller of
+those device functions, such that the kernel attributes are the sum of all
 attributes of all device functions called in this kernel.
 See section 6.7 Attributes for more details.
 
@@ -4353,9 +4353,9 @@ This attribute can be used in both OpenCL and SYCL.
 
 OpenCL documentation:
 The optional attribute intel_reqd_sub_group_size can be used to indicate that
-the kernel must be compiled and executed with the specified subgroup size. When
+the kernel must be compiled and executed with the specified sub group size. When
 this attribute is present, get_max_sub_group_size() is guaranteed to return the
-specified integer value. This is important for the correctness of many subgroup
+specified integer value. This is important for the correctness of many sub group
 algorithms, and in some cases may be used by the compiler to generate more optimal
 code. See `cl_intel_required_subgroup_size
 <https://www.khronos.org/registry/OpenCL/extensions/intel/cl_intel_required_subgroup_size.txt>`
@@ -4370,7 +4370,8 @@ The ``[[intel::sub_group_size(n)]]`` attribute has the same effect as the other
 attribute spellings, except that it follows the SYCL 2020 Attribute Rules. See
  the ``[[intel::named_sub_group_size(NAME)]]`` documentation for clarification.
 
-This attribute is mutually exclusive with ``[[intel::named_sub_group_size(NAME)]]``.
+This attribute is mutually exclusive with ``[[intel::named_sub_group_size(NAME)]]``
+and ``[[intel::sycl_explicit_simd]]``.
 
 In addition to device functions, the required sub-group size attribute may also
 be specified in the definition of a named functor object and lambda functions,
@@ -4400,8 +4401,8 @@ def IntelNamedSubGroupSizeDocs : Documentation {
   let Content = [{
 The ``[[intel::named_sub_group_size(NAME)]]`` attribute works similar to
 ``[[intel::sub_group_size(N)]]`` attribute in that it defines the specific
-subgroup size for the kernel. The ``[[intel::named_sub_group_size(NAME)]]`` form
-accepts a required parameter of either ``automatic`` or ``primary``.
+sub group size for the kernel. The ``[[intel::named_sub_group_size(NAME)]]``
+form accepts a required parameter of either ``automatic`` or ``primary``.
 
 ``automatic`` specifies that the implementation is free to select any of the
 valid sub-group sizes associated with the device to which the kernel is
@@ -4410,7 +4411,8 @@ device's primary sub-group size as reported by
 ``info::device::primary_sub_group_size``.
 
 This attribute may not be combined with ``[[intel::sub_group_size(N)]]``, as
-the two attributes have different meanings.
+the two attributes have different meanings. It is also mutually exclusive with
+``[[intel::sycl_explicit_simd]]``.
 
 In addition to the attributes, a default sub-group size strategy may be
 specified by the ``-fsycl-default-sub-group-size`` command line option, which
@@ -4628,7 +4630,7 @@ Whether a particular pointer may be "null" is an important concern when working
 with pointers in the C family of languages. The various nullability attributes
 indicate whether a particular pointer can be null or not, which makes APIs more
 expressive and can help static analysis tools identify bugs involving null
-pointers. Clang supports several kinds of nullability attributes: the 
+pointers. Clang supports several kinds of nullability attributes: the
 ``nonnull`` and ``returns_nonnull`` attributes indicate which function or
 method parameters and result types can never be null, while nullability type
 qualifiers indicate which pointer types can be null (``_Nullable``) or cannot
@@ -4794,7 +4796,7 @@ memory is not available rather than returning a null pointer:
 The ``returns_nonnull`` attribute implies that returning a null pointer is
 undefined behavior, which the optimizer may take advantage of. The ``_Nonnull``
 type qualifier indicates that a pointer cannot be null in a more general manner
-(because it is part of the type system) and does not imply undefined behavior, 
+(because it is part of the type system) and does not imply undefined behavior,
 making it more widely applicable
 }];
 }
@@ -6615,15 +6617,15 @@ def CFGuardDocs : Documentation {
   let Content = [{
 Code can indicate CFG checks are not wanted with the ``__declspec(guard(nocf))``
 attribute. This directs the compiler to not insert any CFG checks for the entire
-function. This approach is typically used only sparingly in specific situations 
-where the programmer has manually inserted "CFG-equivalent" protection. The 
-programmer knows that they are calling through some read-only function table 
-whose address is obtained through read-only memory references and for which the 
-index is masked to the function table limit. This approach may also be applied 
-to small wrapper functions that are not inlined and that do nothing more than 
-make a call through a function pointer. Since incorrect usage of this directive 
-can compromise the security of CFG, the programmer must be very careful using 
-the directive. Typically, this usage is limited to very small functions that 
+function. This approach is typically used only sparingly in specific situations
+where the programmer has manually inserted "CFG-equivalent" protection. The
+programmer knows that they are calling through some read-only function table
+whose address is obtained through read-only memory references and for which the
+index is masked to the function table limit. This approach may also be applied
+to small wrapper functions that are not inlined and that do nothing more than
+make a call through a function pointer. Since incorrect usage of this directive
+can compromise the security of CFG, the programmer must be very careful using
+the directive. Typically, this usage is limited to very small functions that
 only call one function.
 
 `Control Flow Guard documentation <https://docs.microsoft.com/en-us/windows/win32/secbp/pe-metadata>`
diff --git a/clang/include/clang/Basic/LangOptions.def b/clang/include/clang/Basic/LangOptions.def
index a8a9aa87282aa..d0c80557e3fca 100644
--- a/clang/include/clang/Basic/LangOptions.def
+++ b/clang/include/clang/Basic/LangOptions.def
@@ -425,7 +425,7 @@ LANGOPT(ArmSveVectorBits, 32, 0, "SVE vector size in bits")
 
 ENUM_LANGOPT(DefaultSubGroupSizeType, SubGroupSizeType, 2,
              SubGroupSizeType::None,
-             "Strategy via which sub-group is assigned for SYCL kernel "
+             "Strategy via which sub group is assigned for SYCL kernel "
              "types if not overridden via attributes")
 
 VALUE_LANGOPT(DefaultSubGroupSize, 32, 0,
diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h
index 44d486545ef9f..398dde010c120 100644
--- a/clang/include/clang/Sema/Sema.h
+++ b/clang/include/clang/Sema/Sema.h
@@ -10277,7 +10277,6 @@ class Sema final {
   MergeIntelReqdSubGroupSizeAttr(Decl *D, const IntelReqdSubGroupSizeAttr &A);
   IntelNamedSubGroupSizeAttr *
   MergeIntelNamedSubGroupSizeAttr(Decl *D, const IntelNamedSubGroupSizeAttr &A);
-  SYCLSimdAttr *MergeSYCLSimdAttr(Decl *D, const SYCLSimdAttr &A);
   void AddSYCLIntelNumSimdWorkItemsAttr(Decl *D, const AttributeCommonInfo &CI,
                                         Expr *E);
   SYCLIntelNumSimdWorkItemsAttr *
diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp
index a291693f6e368..7a999bf039cea 100644
--- a/clang/lib/Sema/SemaDecl.cpp
+++ b/clang/lib/Sema/SemaDecl.cpp
@@ -2625,8 +2625,6 @@ static bool mergeDeclAttribute(Sema &S, NamedDecl *D,
     NewAttr = S.MergeIntelReqdSubGroupSizeAttr(D, *A);
   else if (const auto *A = dyn_cast<IntelNamedSubGroupSizeAttr>(Attr))
     NewAttr = S.MergeIntelNamedSubGroupSizeAttr(D, *A);
-  else if (const auto *A = dyn_cast<SYCLSimdAttr>(Attr))
-    NewAttr = S.MergeSYCLSimdAttr(D, *A);
   else if (const auto *A = dyn_cast<SYCLIntelNumSimdWorkItemsAttr>(Attr))
     NewAttr = S.MergeSYCLIntelNumSimdWorkItemsAttr(D, *A);
   else if (const auto *A = dyn_cast<SYCLIntelSchedulerTargetFmaxMhzAttr>(Attr))
diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index 5c9df66894892..bf38982a60545 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -3142,11 +3142,6 @@ static void handleWorkGroupSizeHint(Sema &S, Decl *D, const ParsedAttr &AL) {
 
 void Sema::AddIntelReqdSubGroupSize(Decl *D, const AttributeCommonInfo &CI,
                                     Expr *E) {
-  if (checkAttrMutualExclusion<IntelNamedSubGroupSizeAttr>(*this, D, CI))
-    return;
-  if (checkAttrMutualExclusion<SYCLSimdAttr>(*this, D, CI))
-    return;
-
   if (!E->isValueDependent()) {
     // Validate that we have an integer constant expression and then store the
     // converted constant expression into the semantic attribute so that we
@@ -3187,11 +3182,6 @@ void Sema::AddIntelReqdSubGroupSize(Decl *D, const AttributeCommonInfo &CI,
 IntelReqdSubGroupSizeAttr *
 Sema::MergeIntelReqdSubGroupSizeAttr(Decl *D,
                                      const IntelReqdSubGroupSizeAttr &A) {
-  if (checkAttrMutualExclusion<IntelNamedSubGroupSizeAttr>(*this, D, A))
-    return nullptr;
-  if (checkAttrMutualExclusion<SYCLSimdAttr>(*this, D, A))
-    return nullptr;
-
   // Check to see if there's a duplicate attribute with different values
   // already applied to the declaration.
   if (const auto *DeclAttr = D->getAttr<IntelReqdSubGroupSizeAttr>()) {
@@ -3219,11 +3209,6 @@ static void handleIntelReqdSubGroupSize(Sema &S, Decl *D,
 IntelNamedSubGroupSizeAttr *
 Sema::MergeIntelNamedSubGroupSizeAttr(Decl *D,
                                       const IntelNamedSubGroupSizeAttr &A) {
-  if (checkAttrMutualExclusion<IntelReqdSubGroupSizeAttr>(*this, D, A))
-    return nullptr;
-  if (checkAttrMutualExclusion<SYCLSimdAttr>(*this, D, A))
-    return nullptr;
-
   // Check to see if there's a duplicate attribute with different values
   // already applied to the declaration.
   if (const auto *DeclAttr = D->getAttr<IntelNamedSubGroupSizeAttr>()) {
@@ -3239,11 +3224,6 @@ Sema::MergeIntelNamedSubGroupSizeAttr(Decl *D,
 
 static void handleIntelNamedSubGroupSize(Sema &S, Decl *D,
                                          const ParsedAttr &AL) {
-  if (checkAttrMutualExclusion<IntelReqdSubGroupSizeAttr>(S, D, AL))
-    return;
-  if (checkAttrMutualExclusion<SYCLSimdAttr>(S, D, AL))
-    return;
-
   StringRef SizeStr;
   SourceLocation Loc;
   if (AL.isArgIdent(0)) {
@@ -3262,24 +3242,6 @@ static void handleIntelNamedSubGroupSize(Sema &S, Decl *D,
   D->addAttr(IntelNamedSubGroupSizeAttr::Create(S.Context, SizeType, AL));
 }
 
-SYCLSimdAttr *Sema::MergeSYCLSimdAttr(Decl *D, const SYCLSimdAttr &A) {
-  if (checkAttrMutualExclusion<IntelReqdSubGroupSizeAttr>(*this, D, A))
-    return nullptr;
-  if (checkAttrMutualExclusion<IntelNamedSubGroupSizeAttr>(*this, D, A))
-    return nullptr;
-
-  return A.clone(Context);
-}
-
-static void handleSYCLSimdAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
-  if (checkAttrMutualExclusion<IntelReqdSubGroupSizeAttr>(S, D, AL))
-    return;
-  if (checkAttrMutualExclusion<IntelNamedSubGroupSizeAttr>(S, D, AL))
-    return;
-
-  handleSimpleAttribute<SYCLSimdAttr>(S, D, AL);
-}
-
 void Sema::AddSYCLIntelNumSimdWorkItemsAttr(Decl *D,
                                             const AttributeCommonInfo &CI,
                                             Expr *E) {
@@ -9144,7 +9106,7 @@ static void ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D,
     handleSYCLKernelAttr(S, D, AL);
     break;
   case ParsedAttr::AT_SYCLSimd:
-    handleSYCLSimdAttr(S, D, AL);
+    handleSimpleAttribute<SYCLSimdAttr>(S, D, AL);
     break;
   case ParsedAttr::AT_SYCLDevice:
     handleSYCLDeviceAttr(S, D, AL);
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 87fbeca9abf34..0df9b80b1a0ad 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -681,6 +681,10 @@ class SingleDeviceFunctionTracker {
     return CollectedAttributes;
   }
 
+  llvm::SmallPtrSetImpl<FunctionDecl *> &GetDeviceFunctions() {
+    return DeviceFunctions;
+  }
+
   ~SingleDeviceFunctionTracker() {
     Parent.AddSingleFunction(DeviceFunctions, RecursiveFunctions);
   }
@@ -3795,8 +3799,8 @@ void Sema::MarkDevices() {
     // kernel at a time.
     SingleDeviceFunctionTracker T{Tracker, SYCLKernel};
 
-    CheckSYCL2020Attributes(*this, T.getSYCLKernel(), T.getKernelBody(),
-                            T.getDeviceFunctions());
+    CheckSYCL2020Attributes(*this, T.GetSYCLKernel(), T.GetKernelBody(),
+                            T.GetDeviceFunctions());
     for (auto *A : T.GetCollectedAttributes())
       PropagateAndDiagnoseDeviceAttr(*this, T, A, T.GetSYCLKernel(),
                                      T.GetKernelBody());
diff --git a/clang/test/SemaSYCL/sub-group-size.cpp b/clang/test/SemaSYCL/sub-group-size.cpp
index 2a955def64b5f..414d4b10159e2 100644
--- a/clang/test/SemaSYCL/sub-group-size.cpp
+++ b/clang/test/SemaSYCL/sub-group-size.cpp
@@ -11,14 +11,14 @@
 // expected-note@+1 {{conflicting attribute is here}}
 [[intel::named_sub_group_size(primary)]] [[intel::sub_group_size(1)]] void f2();
 
-// expected-error@+1 {{'sub_group_size' and 'named_sub_group_size' attributes are not compatible}}
-[[intel::sub_group_size(1)]] void f3();
 // expected-note@+1 {{conflicting attribute is here}}
+[[intel::sub_group_size(1)]] void f3();
+// expected-error@+1 {{'named_sub_group_size' and 'sub_group_size' attributes are not compatible}}
 [[intel::named_sub_group_size(primary)]] void f3();
 
-// expected-error@+1 {{'named_sub_group_size' and 'sub_group_size' attributes are not compatible}}
-[[intel::named_sub_group_size(primary)]] void f4();
 // expected-note@+1 {{conflicting attribute is here}}
+[[intel::named_sub_group_size(primary)]] void f4();
+// expected-error@+1 {{'sub_group_size' and 'named_sub_group_size' attributes are not compatible}}
 [[intel::sub_group_size(1)]] void f4();
 
 // expected-note@+1 {{previous attribute is here}}
@@ -41,19 +41,19 @@
 // expected-note@+1 {{conflicting attribute is here}}
 [[intel::sycl_explicit_simd]] [[intel::sub_group_size(1)]] void f9();
 
-// expected-error@+1 {{'named_sub_group_size' and 'sycl_explicit_simd' attributes are not compatible}}
-[[intel::named_sub_group_size(primary)]] void f10();
 // expected-note@+1 {{conflicting attribute is here}}
+[[intel::named_sub_group_size(primary)]] void f10();
+// expected-error@+1 {{'sycl_explicit_simd' and 'named_sub_group_size' attributes are not compatible}}
 [[intel::sycl_explicit_simd]] void f10();
 
-// expected-error@+1 {{'named_sub_group_size' and 'sycl_explicit_simd' attributes are not compatible}}
-[[intel::named_sub_group_size("primary")]] void f11();
 // expected-note@+1 {{conflicting attribute is here}}
+[[intel::named_sub_group_size("primary")]] void f11();
+// expected-error@+1 {{'sycl_explicit_simd' and 'named_sub_group_size' attributes are not compatible}}
 [[intel::sycl_explicit_simd]] void f11();
 
-// expected-error@+1 {{'named_sub_group_size' and 'sycl_explicit_simd' attributes are not compatible}}
-[[intel::named_sub_group_size("automatic")]] void f12();
 // expected-note@+1 {{conflicting attribute is here}}
+[[intel::named_sub_group_size("automatic")]] void f12();
+// expected-error@+1 {{'sycl_explicit_simd' and 'named_sub_group_size' attributes are not compatible}}
 [[intel::sycl_explicit_simd]] void f12();
 
 // expected-warning@+1 {{'named_sub_group_size' attribute argument not supported: invalid string}}

From f0281660e4b01d51daa32031d3a41f9f7257d944 Mon Sep 17 00:00:00 2001
From: Erich Keane <erich.keane@intel.com>
Date: Mon, 12 Apr 2021 10:24:25 -0700
Subject: [PATCH 15/22] Revert some inadvertent changes to AttrDocs.td

---
 clang/include/clang/Basic/AttrDocs.td | 36 +++++++++++++--------------
 1 file changed, 18 insertions(+), 18 deletions(-)

diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td
index 5978c7a4573f2..f7a75b9fdd0da 100644
--- a/clang/include/clang/Basic/AttrDocs.td
+++ b/clang/include/clang/Basic/AttrDocs.td
@@ -453,16 +453,16 @@ def NoMergeDocs : Documentation {
   let Category = DocCatStmt;
   let Content = [{
 If a statement is marked ``nomerge`` and contains call expressions, those call
-expressions inside the statement will not be merged during optimization. This
+expressions inside the statement will not be merged during optimization. This 
 attribute can be used to prevent the optimizer from obscuring the source
 location of certain calls. For example, it will prevent tail merging otherwise
 identical code sequences that raise an exception or terminate the program. Tail
 merging normally reduces the precision of source location information, making
 stack traces less useful for debugging. This attribute gives the user control
-over the tradeoff between code size and debug information precision.
+over the tradeoff between code size and debug information precision. 
 
-``nomerge`` attribute can also be used as function attribute to prevent all
-calls to the specified function from merging. It has no effect on indirect
+``nomerge`` attribute can also be used as function attribute to prevent all 
+calls to the specified function from merging. It has no effect on indirect 
 calls.
   }];
 }
@@ -2544,11 +2544,11 @@ argument to **clEnqueueNDRangeKernel** (in OpenCL) or to
 generated code appropriately for the kernel to which attribute is applied.
 
 While semantic of this attribute is the same between OpenCL and SYCL,
-spelling is a bit different:
+spelling is a bit different: 
 
 SYCL 1.2.1 describes ``[[cl::reqd_work_group_size(X, Y, Z)]]`` spelling: this
-attribute is legal on device functions and is propagated down to any caller of
-those device functions, such that the kernel attributes are the sum of all
+attribute is legal on device functions and is propagated down to any caller of 
+those device functions, such that the kernel attributes are the sum of all 
 attributes of all device functions called in this kernel.
 See section 6.7 Attributes for more details.
 
@@ -4630,7 +4630,7 @@ Whether a particular pointer may be "null" is an important concern when working
 with pointers in the C family of languages. The various nullability attributes
 indicate whether a particular pointer can be null or not, which makes APIs more
 expressive and can help static analysis tools identify bugs involving null
-pointers. Clang supports several kinds of nullability attributes: the
+pointers. Clang supports several kinds of nullability attributes: the 
 ``nonnull`` and ``returns_nonnull`` attributes indicate which function or
 method parameters and result types can never be null, while nullability type
 qualifiers indicate which pointer types can be null (``_Nullable``) or cannot
@@ -4796,7 +4796,7 @@ memory is not available rather than returning a null pointer:
 The ``returns_nonnull`` attribute implies that returning a null pointer is
 undefined behavior, which the optimizer may take advantage of. The ``_Nonnull``
 type qualifier indicates that a pointer cannot be null in a more general manner
-(because it is part of the type system) and does not imply undefined behavior,
+(because it is part of the type system) and does not imply undefined behavior, 
 making it more widely applicable
 }];
 }
@@ -6617,15 +6617,15 @@ def CFGuardDocs : Documentation {
   let Content = [{
 Code can indicate CFG checks are not wanted with the ``__declspec(guard(nocf))``
 attribute. This directs the compiler to not insert any CFG checks for the entire
-function. This approach is typically used only sparingly in specific situations
-where the programmer has manually inserted "CFG-equivalent" protection. The
-programmer knows that they are calling through some read-only function table
-whose address is obtained through read-only memory references and for which the
-index is masked to the function table limit. This approach may also be applied
-to small wrapper functions that are not inlined and that do nothing more than
-make a call through a function pointer. Since incorrect usage of this directive
-can compromise the security of CFG, the programmer must be very careful using
-the directive. Typically, this usage is limited to very small functions that
+function. This approach is typically used only sparingly in specific situations 
+where the programmer has manually inserted "CFG-equivalent" protection. The 
+programmer knows that they are calling through some read-only function table 
+whose address is obtained through read-only memory references and for which the 
+index is masked to the function table limit. This approach may also be applied 
+to small wrapper functions that are not inlined and that do nothing more than 
+make a call through a function pointer. Since incorrect usage of this directive 
+can compromise the security of CFG, the programmer must be very careful using 
+the directive. Typically, this usage is limited to very small functions that 
 only call one function.
 
 `Control Flow Guard documentation <https://docs.microsoft.com/en-us/windows/win32/secbp/pe-metadata>`

From 1cd4cc26c13636e36c81f03c878313d1dafb6a62 Mon Sep 17 00:00:00 2001
From: Erich Keane <erich.keane@intel.com>
Date: Mon, 12 Apr 2021 12:05:32 -0700
Subject: [PATCH 16/22] Clang-format fixes + getting codegen to work right

---
 clang/lib/Frontend/CompilerInvocation.cpp |  1 -
 clang/lib/Sema/SemaSYCL.cpp               | 32 ++++++++++++++++-------
 clang/test/CodeGenSYCL/sub-group-size.cpp |  2 +-
 clang/test/SemaSYCL/kernel-handler.cpp    |  3 +++
 clang/test/SemaSYCL/wrapped-accessor.cpp  |  1 +
 5 files changed, 28 insertions(+), 11 deletions(-)

diff --git a/clang/lib/Frontend/CompilerInvocation.cpp b/clang/lib/Frontend/CompilerInvocation.cpp
index 0d615d07568c3..e6b74741c713b 100644
--- a/clang/lib/Frontend/CompilerInvocation.cpp
+++ b/clang/lib/Frontend/CompilerInvocation.cpp
@@ -3516,7 +3516,6 @@ void CompilerInvocation::GenerateLangArgs(const LangOptions &Opts,
   case LangOptions::SubGroupSizeType::None:
     break;
   }
-
 }
 
 bool CompilerInvocation::ParseLangArgs(LangOptions &Opts, ArgList &Args,
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 0df9b80b1a0ad..f49afce5cca64 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -558,8 +558,6 @@ class DeviceFunctionTracker {
       if (const FunctionDecl *Def = FD->getDefinition())
         Diagnoser.CheckBody(Def->getBody());
   }
-
-
 };
 
 // This type does the heavy lifting for the management of device functions,
@@ -2297,6 +2295,14 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
 
     BodyStmts.insert(BodyStmts.end(), FinalizeStmts.begin(),
                      FinalizeStmts.end());
+
+    // Make sure that this is marked as a kernel so that the code-gen can make
+    // decisions based on that. We cannot add this earlier, otherwise the call
+    // to TransformStmt in replaceWithLocalClone can diagnose something that got
+    // diagnosed on the actual kernel.
+    KernelObjClone->addAttr(
+        SYCLKernelAttr::CreateImplicit(SemaRef.getASTContext()));
+
     return CompoundStmt::Create(SemaRef.getASTContext(), BodyStmts, {}, {});
   }
 
@@ -3636,10 +3642,19 @@ static void CheckSYCL2020Attributes(
     Sema &S, FunctionDecl *SYCLKernel, FunctionDecl *KernelBody,
     const llvm::SmallPtrSetImpl<FunctionDecl *> &CalledFuncs) {
 
-  // If the kernel has a body, we should get the attributes for the kernel from
-  // there instead, so that we get the functor object.
-  if (KernelBody)
+  if (KernelBody) {
+    // Make sure the kernel itself has all the 2020 attributes, since we don't
+    // do propagation of these.
+    if (auto *A = KernelBody->getAttr<IntelReqdSubGroupSizeAttr>())
+      if (A->isSYCL2020Spelling())
+        SYCLKernel->addAttr(A);
+    if (auto *A = KernelBody->getAttr<IntelNamedSubGroupSizeAttr>())
+      SYCLKernel->addAttr(A);
+
+    // If the kernel has a body, we should get the attributes for the kernel from
+    // there instead, so that we get the functor object.
     SYCLKernel = KernelBody;
+  }
 
   for (auto *FD : CalledFuncs) {
     if (FD == SYCLKernel || FD == KernelBody)
@@ -3647,10 +3662,9 @@ static void CheckSYCL2020Attributes(
     for (auto *Attr : FD->attrs()) {
       switch (Attr->getKind()) {
       case attr::Kind::IntelReqdSubGroupSize:
-        if (const auto *A = cast<IntelReqdSubGroupSizeAttr>(Attr))
-          // Pre SYCL2020 spellings handled during collection.
-          if (!A->isSYCL2020Spelling())
-            break;
+        // Pre SYCL2020 spellings handled during collection.
+        if (!cast<IntelReqdSubGroupSizeAttr>(Attr)->isSYCL2020Spelling())
+          break;
         LLVM_FALLTHROUGH;
       case attr::Kind::IntelNamedSubGroupSize:
         CheckSYCL2020SubGroupSizes(S, SYCLKernel, FD);
diff --git a/clang/test/CodeGenSYCL/sub-group-size.cpp b/clang/test/CodeGenSYCL/sub-group-size.cpp
index 911566fb4a439..4626ab3e5516e 100644
--- a/clang/test/CodeGenSYCL/sub-group-size.cpp
+++ b/clang/test/CodeGenSYCL/sub-group-size.cpp
@@ -6,6 +6,6 @@
 using namespace cl::sycl;
 
 void default_behavior() {
-  kernel_single_task<class Kernel1>([]() {
+  kernel_single_task<class Kernel1>([]() /*[[intel::reqd_sub_group_size(10)]]*/ {
   });
 }
diff --git a/clang/test/SemaSYCL/kernel-handler.cpp b/clang/test/SemaSYCL/kernel-handler.cpp
index a5df0e186e10b..087c8c0c66dd3 100644
--- a/clang/test/SemaSYCL/kernel-handler.cpp
+++ b/clang/test/SemaSYCL/kernel-handler.cpp
@@ -39,6 +39,7 @@ int main() {
 // NONATIVESUPPORT-NEXT: InitListExpr
 // NONATIVESUPPORT-NEXT: ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
 // NONATIVESUPPORT-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_' 'int'
+// NONATIVESUPPORT-NEXT: SYCLKernelAttr
 
 // Check declaration and initialization of kernel handler local clone using default constructor
 // NONATIVESUPPORT-NEXT: DeclStmt
@@ -74,6 +75,7 @@ int main() {
 // NONATIVESUPPORT-NEXT: InitListExpr
 // NONATIVESUPPORT-NEXT: ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
 // NONATIVESUPPORT-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_' 'int'
+// NONATIVESUPPORT-NEXT: SYCLKernelAttr
 
 // Check declaration and initialization of kernel handler local clone using default constructor
 // NONATIVESUPPORT-NEXT: DeclStmt
@@ -115,6 +117,7 @@ int main() {
 // NATIVESUPPORT-NEXT: InitListExpr
 // NATIVESUPPORT-NEXT: ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
 // NATIVESUPPORT-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_' 'int'
+// NATIVESUPPORT-NEXT: SYCLKernelAttr
 
 // Check declaration and initialization of kernel handler local clone using default constructor
 // NATIVESUPPORT-NEXT: DeclStmt
diff --git a/clang/test/SemaSYCL/wrapped-accessor.cpp b/clang/test/SemaSYCL/wrapped-accessor.cpp
index b8d79501c1fda..855c6ad561486 100644
--- a/clang/test/SemaSYCL/wrapped-accessor.cpp
+++ b/clang/test/SemaSYCL/wrapped-accessor.cpp
@@ -39,6 +39,7 @@ int main() {
 // CHECK-NEXT: InitListExpr {{.*}}'(lambda at {{.*}}wrapped-accessor.cpp{{.*}})'
 // CHECK-NEXT: InitListExpr {{.*}}'AccWrapper<sycl::accessor<int, 1, sycl::access::mode::read_write, sycl::access::target::global_buffer, sycl::access::placeholder::false_t>>'
 // CHECK-NEXT: CXXConstructExpr {{.*}}'sycl::accessor<int, 1, sycl::access::mode::read_write, sycl::access::target::global_buffer, sycl::access::placeholder::false_t>':'sycl::accessor<int, 1, sycl::access::mode::read_write, sycl::access::target::global_buffer, sycl::access::placeholder::false_t>' 'void () noexcept'
+// CHECK-NEXT: SYCLKernelAttr
 
 // Check that accessor field of the wrapper object is initialized using __init method
 // CHECK-NEXT: CXXMemberCallExpr {{.*}} 'void'

From c1ab6a0d3ccbe3133fc716a2d200038002ff8ff3 Mon Sep 17 00:00:00 2001
From: Erich Keane <erich.keane@intel.com>
Date: Tue, 13 Apr 2021 06:50:05 -0700
Subject: [PATCH 17/22] Add sufficient code-gen tests, fix codegen.

I had some other issues getting code gen set up, since the esimd and the
way sycl-device metadata was generated, so I had to fix that along the
way!
---
 clang/lib/CodeGen/CodeGenFunction.cpp         |  5 +--
 clang/lib/Sema/SemaSYCL.cpp                   | 18 +++++-----
 clang/test/CodeGenSYCL/esimd_metadata2.cpp    |  4 +--
 clang/test/CodeGenSYCL/sub-group-size.cpp     | 35 ++++++++++++++++---
 .../SemaSYCL/allow-constexpr-recursion.cpp    |  4 +--
 clang/test/SemaSYCL/kernel-handler.cpp        |  3 --
 clang/test/SemaSYCL/wrapped-accessor.cpp      |  1 -
 7 files changed, 47 insertions(+), 23 deletions(-)

diff --git a/clang/lib/CodeGen/CodeGenFunction.cpp b/clang/lib/CodeGen/CodeGenFunction.cpp
index 1c09702a07504..47798e62c5404 100644
--- a/clang/lib/CodeGen/CodeGenFunction.cpp
+++ b/clang/lib/CodeGen/CodeGenFunction.cpp
@@ -592,7 +592,7 @@ CodeGenFunction::DecodeAddrUsedInPrologue(llvm::Value *F,
 void CodeGenFunction::EmitOpenCLKernelMetadata(const FunctionDecl *FD,
                                                llvm::Function *Fn)
 {
-  if (!FD->hasAttr<OpenCLKernelAttr>())
+  if (!FD->hasAttr<OpenCLKernelAttr>() && !FD->hasAttr<SYCLDeviceAttr>())
     return;
 
   // TODO Module identifier is not reliable for this purpose since two modules
@@ -602,7 +602,8 @@ void CodeGenFunction::EmitOpenCLKernelMetadata(const FunctionDecl *FD,
 
   llvm::LLVMContext &Context = getLLVMContext();
 
-  CGM.GenOpenCLArgMetadata(Fn, FD, this);
+  if (FD->hasAttr<OpenCLKernelAttr>())
+    CGM.GenOpenCLArgMetadata(Fn, FD, this);
 
   if (const VecTypeHintAttr *A = FD->getAttr<VecTypeHintAttr>()) {
     QualType HintQTy = A->getTypeHint();
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index f49afce5cca64..a682d56bd7eb8 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -1937,6 +1937,13 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
     KernelDecl->setType(FuncType);
     KernelDecl->setParams(Params);
 
+    // Make sure that this is marked as a kernel so that the code-gen can make
+    // decisions based on that. We cannot add this earlier, otherwise the call
+    // to TransformStmt in replaceWithLocalClone can diagnose something that got
+    // diagnosed on the actual kernel.
+    KernelDecl->addAttr(
+        SYCLKernelAttr::CreateImplicit(SemaRef.getASTContext()));
+
     SemaRef.addSyclDeviceDecl(KernelDecl);
   }
 
@@ -2296,13 +2303,6 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
     BodyStmts.insert(BodyStmts.end(), FinalizeStmts.begin(),
                      FinalizeStmts.end());
 
-    // Make sure that this is marked as a kernel so that the code-gen can make
-    // decisions based on that. We cannot add this earlier, otherwise the call
-    // to TransformStmt in replaceWithLocalClone can diagnose something that got
-    // diagnosed on the actual kernel.
-    KernelObjClone->addAttr(
-        SYCLKernelAttr::CreateImplicit(SemaRef.getASTContext()));
-
     return CompoundStmt::Create(SemaRef.getASTContext(), BodyStmts, {}, {});
   }
 
@@ -3651,8 +3651,8 @@ static void CheckSYCL2020Attributes(
     if (auto *A = KernelBody->getAttr<IntelNamedSubGroupSizeAttr>())
       SYCLKernel->addAttr(A);
 
-    // If the kernel has a body, we should get the attributes for the kernel from
-    // there instead, so that we get the functor object.
+    // If the kernel has a body, we should get the attributes for the kernel
+    // from there instead, so that we get the functor object.
     SYCLKernel = KernelBody;
   }
 
diff --git a/clang/test/CodeGenSYCL/esimd_metadata2.cpp b/clang/test/CodeGenSYCL/esimd_metadata2.cpp
index 5c8ebb09f26b0..3d3d53920a798 100644
--- a/clang/test/CodeGenSYCL/esimd_metadata2.cpp
+++ b/clang/test/CodeGenSYCL/esimd_metadata2.cpp
@@ -8,8 +8,8 @@ __attribute__((sycl_device)) void shared_func() { shared_func_decl(); }
 
 __attribute__((sycl_device)) __attribute__((sycl_explicit_simd)) void esimd_func() { shared_func(); }
 
-// CHECK-ESIMD-DAG: define {{.*}}spir_kernel void @{{.*}}kernel_cm() #{{[0-9]+}} !sycl_explicit_simd !{{[0-9]+}} {{.*}} !intel_reqd_sub_group_size ![[SGSIZE1:[0-9]+]] {{.*}}{
-// CHECK-ESIMD-DAG: define {{.*}}spir_func void @{{.*}}esimd_funcv() #{{[0-9]+}} !sycl_explicit_simd !{{[0-9]+}} {
+// CHECK-ESIMD-DAG: define {{.*}}spir_kernel void @{{.*}}kernel_cm() #{{[0-9]+}} !sycl_explicit_simd !{{[0-9]+}} !intel_reqd_sub_group_size ![[SGSIZE1:[0-9]+]] {{.*}}{
+// CHECK-ESIMD-DAG: define {{.*}}spir_func void @{{.*}}esimd_funcv() #{{[0-9]+}} !sycl_explicit_simd !{{[0-9]+}} !intel_reqd_sub_group_size ![[SGSIZE1]] {
 // CHECK-ESIMD-DAG: define {{.*}}spir_func void @{{.*}}shared_funcv() #{{[0-9]+}} {
 // CHECK-ESIMD-DAG: define linkonce_odr spir_func void @_ZN12ESIMDFunctorclEv({{.*}}) #{{[0-9]+}} {{.*}} !sycl_explicit_simd !{{[0-9]+}} {
 // CHECK-ESIMD-DAG: declare spir_func void @{{.*}}shared_func_declv() #{{[0-9]+}}
diff --git a/clang/test/CodeGenSYCL/sub-group-size.cpp b/clang/test/CodeGenSYCL/sub-group-size.cpp
index 4626ab3e5516e..2f7be0d4c1d20 100644
--- a/clang/test/CodeGenSYCL/sub-group-size.cpp
+++ b/clang/test/CodeGenSYCL/sub-group-size.cpp
@@ -1,11 +1,38 @@
-// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2020 -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s --check-prefix=NONE
-// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -fsycl-default-sub-group-size=primary -sycl-std=2020 -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s --check-prefix=PRIM_DEF
-// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -fsycl-default-sub-group-size=10 -sycl-std=2020 -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s --check-prefix=TEN_DEF
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2020 -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s --check-prefixes=NONE,ALL
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -fsycl-default-sub-group-size=primary -sycl-std=2020 -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s --check-prefixes=PRIM_DEF,ALL
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -fsycl-default-sub-group-size=10 -sycl-std=2020 -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s --check-prefixes=TEN_DEF,ALL
 
 #include "Inputs/sycl.hpp"
 using namespace cl::sycl;
 
+[[intel::named_sub_group_size(primary)]] SYCL_EXTERNAL void external_primary() {}
+// ALL-DAG: define {{.*}}spir_func void @{{.*}}external_primary{{.*}}() #{{[0-9]+}} !intel_reqd_sub_group_size ![[PRIMARY:[0-9]+]]
+
+[[intel::sub_group_size(10)]] SYCL_EXTERNAL void external_10() {}
+// ALL-DAG: define {{.*}}spir_func void @{{.*}}external_10{{.*}}() #{{[0-9]+}} !intel_reqd_sub_group_size ![[TEN:[0-9]+]]
+
+SYCL_EXTERNAL void external_default_behavior() {}
+// NONE-DAG: define {{.*}}spir_func void @{{.*}}external_default_behavior{{.*}}() #{{[0-9]+}} {
+// PRIM_DEF-DAG: define {{.*}}spir_func void @{{.*}}external_default_behavior{{.*}}() #{{[0-9]+}} !intel_reqd_sub_group_size ![[PRIMARY]] {
+// TEN_DEF-DAG: define {{.*}}spir_func void @{{.*}}external_default_behavior{{.*}}() #{{[0-9]+}} !intel_reqd_sub_group_size ![[TEN]] {
+
 void default_behavior() {
-  kernel_single_task<class Kernel1>([]() /*[[intel::reqd_sub_group_size(10)]]*/ {
+  kernel_single_task<class Kernel1>([]() {
   });
 }
+// NONE-DAG: define {{.*}}spir_kernel void @{{.*}}Kernel1() #{{[0-9]+}} !kernel_arg_buffer_location !{{[0-9]+}} {
+// PRIM_DEF-DAG: define {{.*}}spir_kernel void @{{.*}}Kernel1() #{{[0-9]+}} !intel_reqd_sub_group_size ![[PRIMARY]]
+// TEN_DEF-DAG: define {{.*}}spir_kernel void @{{.*}}Kernel1() #{{[0-9]+}} !intel_reqd_sub_group_size ![[TEN]]
+
+void primary() {
+  kernel_single_task<class Kernel2>([]() [[intel::named_sub_group_size(primary)]]{});
+}
+// ALL-DAG: define {{.*}}spir_kernel void @{{.*}}Kernel2() #{{[0-9]+}} !intel_reqd_sub_group_size ![[PRIMARY]]
+
+void ten() {
+  kernel_single_task<class Kernel3>([]() [[intel::sub_group_size(10)]]{});
+}
+// ALL-DAG: define {{.*}}spir_kernel void @{{.*}}Kernel3() #{{[0-9]+}} !intel_reqd_sub_group_size ![[TEN]]
+
+// PRIM_DEF: ![[PRIMARY]] = !{!"primary"}
+// TEN_DEF: ![[TEN]] = !{i32 10}
diff --git a/clang/test/SemaSYCL/allow-constexpr-recursion.cpp b/clang/test/SemaSYCL/allow-constexpr-recursion.cpp
index a73ff61280ce6..4dde146900849 100644
--- a/clang/test/SemaSYCL/allow-constexpr-recursion.cpp
+++ b/clang/test/SemaSYCL/allow-constexpr-recursion.cpp
@@ -36,7 +36,7 @@ struct ConditionallyExplicitCtor {
 
 void conditionally_noexcept() noexcept(constexpr_recurse(5)) {}
 
-template<int I>
+template <int I>
 void ConstexprIf1() {
   if constexpr (I == 1)
     ConstexprIf1<I>();
@@ -44,7 +44,7 @@ void ConstexprIf1() {
 
 // Same as the above, but split up so the diagnostic is more clear.
 // expected-note@+2 2{{function implemented using recursion declared here}}
-template<int I>
+template <int I>
 void ConstexprIf2() {
   if constexpr (I == 1)
     // expected-error@+1{{SYCL kernel cannot call a recursive function}}
diff --git a/clang/test/SemaSYCL/kernel-handler.cpp b/clang/test/SemaSYCL/kernel-handler.cpp
index 087c8c0c66dd3..a5df0e186e10b 100644
--- a/clang/test/SemaSYCL/kernel-handler.cpp
+++ b/clang/test/SemaSYCL/kernel-handler.cpp
@@ -39,7 +39,6 @@ int main() {
 // NONATIVESUPPORT-NEXT: InitListExpr
 // NONATIVESUPPORT-NEXT: ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
 // NONATIVESUPPORT-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_' 'int'
-// NONATIVESUPPORT-NEXT: SYCLKernelAttr
 
 // Check declaration and initialization of kernel handler local clone using default constructor
 // NONATIVESUPPORT-NEXT: DeclStmt
@@ -75,7 +74,6 @@ int main() {
 // NONATIVESUPPORT-NEXT: InitListExpr
 // NONATIVESUPPORT-NEXT: ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
 // NONATIVESUPPORT-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_' 'int'
-// NONATIVESUPPORT-NEXT: SYCLKernelAttr
 
 // Check declaration and initialization of kernel handler local clone using default constructor
 // NONATIVESUPPORT-NEXT: DeclStmt
@@ -117,7 +115,6 @@ int main() {
 // NATIVESUPPORT-NEXT: InitListExpr
 // NATIVESUPPORT-NEXT: ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
 // NATIVESUPPORT-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_' 'int'
-// NATIVESUPPORT-NEXT: SYCLKernelAttr
 
 // Check declaration and initialization of kernel handler local clone using default constructor
 // NATIVESUPPORT-NEXT: DeclStmt
diff --git a/clang/test/SemaSYCL/wrapped-accessor.cpp b/clang/test/SemaSYCL/wrapped-accessor.cpp
index 855c6ad561486..b8d79501c1fda 100644
--- a/clang/test/SemaSYCL/wrapped-accessor.cpp
+++ b/clang/test/SemaSYCL/wrapped-accessor.cpp
@@ -39,7 +39,6 @@ int main() {
 // CHECK-NEXT: InitListExpr {{.*}}'(lambda at {{.*}}wrapped-accessor.cpp{{.*}})'
 // CHECK-NEXT: InitListExpr {{.*}}'AccWrapper<sycl::accessor<int, 1, sycl::access::mode::read_write, sycl::access::target::global_buffer, sycl::access::placeholder::false_t>>'
 // CHECK-NEXT: CXXConstructExpr {{.*}}'sycl::accessor<int, 1, sycl::access::mode::read_write, sycl::access::target::global_buffer, sycl::access::placeholder::false_t>':'sycl::accessor<int, 1, sycl::access::mode::read_write, sycl::access::target::global_buffer, sycl::access::placeholder::false_t>' 'void () noexcept'
-// CHECK-NEXT: SYCLKernelAttr
 
 // Check that accessor field of the wrapper object is initialized using __init method
 // CHECK-NEXT: CXXMemberCallExpr {{.*}} 'void'

From 9e186cf6aac68a5406eaf9fb49c0f4c13bc11698 Mon Sep 17 00:00:00 2001
From: Erich Keane <erich.keane@intel.com>
Date: Tue, 13 Apr 2021 10:56:40 -0700
Subject: [PATCH 18/22] Fix Aaron's nits

---
 clang/include/clang/Basic/Attr.td     | 4 ++--
 clang/lib/CodeGen/CodeGenFunction.cpp | 3 +--
 clang/lib/Sema/SemaSYCL.cpp           | 6 +++---
 3 files changed, 6 insertions(+), 7 deletions(-)

diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td
index 04fd3e64ea744..afdbd8ce93bc2 100644
--- a/clang/include/clang/Basic/Attr.td
+++ b/clang/include/clang/Basic/Attr.td
@@ -1455,8 +1455,8 @@ def IntelNamedSubGroupSize : InheritableAttr {
   let SupportsNonconformingLambdaSyntax = 1;
 }
 
-def : MutualExclusions<
-          [IntelReqdSubGroupSize, IntelNamedSubGroupSize, SYCLSimd]>;
+def :
+  MutualExclusions<[IntelReqdSubGroupSize, IntelNamedSubGroupSize, SYCLSimd]>;
 
 // This attribute is both a type attribute, and a declaration attribute (for
 // parameter variables).
diff --git a/clang/lib/CodeGen/CodeGenFunction.cpp b/clang/lib/CodeGen/CodeGenFunction.cpp
index 47798e62c5404..223f5a9fdb887 100644
--- a/clang/lib/CodeGen/CodeGenFunction.cpp
+++ b/clang/lib/CodeGen/CodeGenFunction.cpp
@@ -676,8 +676,7 @@ void CodeGenFunction::EmitOpenCLKernelMetadata(const FunctionDecl *FD,
   // SCYL2020 doesn't propagate attributes, so don't put it in an intermediate
   // location.
   if (IsKernelOrDevice) {
-    if (const IntelNamedSubGroupSizeAttr *A =
-            FD->getAttr<IntelNamedSubGroupSizeAttr>()) {
+    if (const auto *A = FD->getAttr<IntelNamedSubGroupSizeAttr>()) {
       llvm::Metadata *AttrMDArgs[] = {llvm::MDString::get(
           Context, A->getType() == IntelNamedSubGroupSizeAttr::Primary
                        ? "primary"
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index a682d56bd7eb8..d8dcff0149421 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -3574,7 +3574,7 @@ void Sema::ConstructOpenCLKernel(FunctionDecl *KernelCallerFunc,
 // attributes, then the global settings.
 static std::pair<LangOptions::SubGroupSizeType, int64_t>
 CalcEffectiveSubGroup(ASTContext &Ctx, const LangOptions &LO,
-                      FunctionDecl *FD) {
+                      const FunctionDecl *FD) {
   if (const auto *A = FD->getAttr<IntelReqdSubGroupSizeAttr>()) {
     int64_t Val = getIntExprValue(A->getValue(), Ctx);
     return {LangOptions::SubGroupSizeType::Integer, Val};
@@ -3591,7 +3591,7 @@ CalcEffectiveSubGroup(ASTContext &Ctx, const LangOptions &LO,
           static_cast<uint64_t>(LO.DefaultSubGroupSize)};
 }
 
-static SourceLocation GetSubGroupLoc(FunctionDecl *FD) {
+static SourceLocation GetSubGroupLoc(const FunctionDecl *FD) {
   if (const auto *A = FD->getAttr<IntelReqdSubGroupSizeAttr>())
     return A->getLocation();
   if (const auto *A = FD->getAttr<IntelNamedSubGroupSizeAttr>())
@@ -3600,7 +3600,7 @@ static SourceLocation GetSubGroupLoc(FunctionDecl *FD) {
 }
 
 static void CheckSYCL2020SubGroupSizes(Sema &S, FunctionDecl *SYCLKernel,
-                                       FunctionDecl *FD) {
+                                       const FunctionDecl *FD) {
   // If they are the same, no error.
   if (CalcEffectiveSubGroup(S.Context, S.getLangOpts(), SYCLKernel) ==
       CalcEffectiveSubGroup(S.Context, S.getLangOpts(), FD))

From 59b7b5b0881695c84cc29eedda1c1b4461c06101 Mon Sep 17 00:00:00 2001
From: Erich Keane <erich.keane@intel.com>
Date: Tue, 13 Apr 2021 11:43:41 -0700
Subject: [PATCH 19/22] Add driver test, fix it the way Mike suggested

---
 clang/lib/Driver/ToolChains/Clang.cpp     | 6 +++---
 clang/test/Driver/sycl-sub-group-size.cpp | 8 ++++++++
 2 files changed, 11 insertions(+), 3 deletions(-)
 create mode 100644 clang/test/Driver/sycl-sub-group-size.cpp

diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp
index 683a66bbbcdf1..9973f74b0cc1d 100644
--- a/clang/lib/Driver/ToolChains/Clang.cpp
+++ b/clang/lib/Driver/ToolChains/Clang.cpp
@@ -5962,9 +5962,6 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA,
   if (Args.hasFlag(options::OPT_fsycl, options::OPT_fno_sycl, false))
     Args.AddLastArg(CmdArgs, options::OPT_sycl_std_EQ);
 
-  // Forward -fsycl-default-sub-group-size if in SYCL mode.
-  if (Args.hasFlag(options::OPT_fsycl, options::OPT_fno_sycl, false))
-    Args.AddLastArg(CmdArgs, options::OPT_fsycl_default_sub_group_size);
 
   // Forward -fsycl-instrument-device-code option to cc1. This option can only
   // be used with spir triple.
@@ -6682,6 +6679,9 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA,
         HeaderOpt.append(Header);
         CmdArgs.push_back(Args.MakeArgString(HeaderOpt));
       }
+
+      // Forward -fsycl-default-sub-group-size if in SYCL mode.
+      Args.AddLastArg(CmdArgs, options::OPT_fsycl_default_sub_group_size);
     }
 
     if (Args.hasArg(options::OPT_fsycl_unnamed_lambda))
diff --git a/clang/test/Driver/sycl-sub-group-size.cpp b/clang/test/Driver/sycl-sub-group-size.cpp
new file mode 100644
index 0000000000000..151ba82fc873e
--- /dev/null
+++ b/clang/test/Driver/sycl-sub-group-size.cpp
@@ -0,0 +1,8 @@
+// RUN: %clang -### -fsycl %s 2>&1 | FileCheck %s -check-prefix=CHECK-DEFAULT
+// CHECK-DEFAULT-NOT: "fsycl-default-sub-group-size"
+
+// RUN: %clang -### -fsycl -fsycl-default-sub-group-size=primary %s 2>&1 | FileCheck %s -check-prefix=PRIM
+// PRIM: "-fsycl-default-sub-group-size" "primary"
+
+// RUN: %clang -### -fsycl -fsycl-default-sub-group-size=10 %s 2>&1 | FileCheck %s -check-prefix=TEN
+// TEN: "-fsycl-default-sub-group-size" "10"

From 3129b5f84fb97ff67dccb1b649ecc1e9ffb74ae8 Mon Sep 17 00:00:00 2001
From: Erich Keane <erich.keane@intel.com>
Date: Tue, 13 Apr 2021 11:48:07 -0700
Subject: [PATCH 20/22] Remove newline to satisfy clang-format

---
 clang/lib/Driver/ToolChains/Clang.cpp | 1 -
 1 file changed, 1 deletion(-)

diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp
index 9973f74b0cc1d..5d945fa9edb4d 100644
--- a/clang/lib/Driver/ToolChains/Clang.cpp
+++ b/clang/lib/Driver/ToolChains/Clang.cpp
@@ -5962,7 +5962,6 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA,
   if (Args.hasFlag(options::OPT_fsycl, options::OPT_fno_sycl, false))
     Args.AddLastArg(CmdArgs, options::OPT_sycl_std_EQ);
 
-
   // Forward -fsycl-instrument-device-code option to cc1. This option can only
   // be used with spir triple.
   if (Arg *A = Args.getLastArg(options::OPT_fsycl_instrument_device_code)) {

From 4fc91bd43625956e2990ce2904b3271fa2c3d428 Mon Sep 17 00:00:00 2001
From: Erich Keane <erich.keane@intel.com>
Date: Tue, 13 Apr 2021 13:08:27 -0700
Subject: [PATCH 21/22] Fix check-sycl tests by excluding old-spelling from the
 test

Really we should probably be checking this, but since it is a pattern
that is apparently commonly used, AND is a breaking change, we won't
enforce this rule. The standard doesn't require us to (since it is
pretty silent on mixing attribute types).
---
 clang/lib/Sema/SemaSYCL.cpp            | 7 +++++++
 clang/test/SemaSYCL/sub-group-size.cpp | 7 +++++++
 2 files changed, 14 insertions(+)

diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index d8dcff0149421..e068c9021965a 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -3626,6 +3626,13 @@ static void CheckSYCL2020SubGroupSizes(Sema &S, FunctionDecl *SYCLKernel,
 
   // Else this doesn't have an attribute, which can only be caused by this being
   // an undefined SYCL_EXTERNAL, and the kernel has an attribute that conflicts.
+  if (const auto *A = SYCLKernel->getAttr<IntelReqdSubGroupSizeAttr>()) {
+    // Don't diagnose this if the kernel got its size from the 'old' attribute
+    // spelling.
+    if (!A->isSYCL2020Spelling())
+      return;
+  }
+
   assert(KernelAttrLoc.isValid() && "Kernel doesn't have attribute either?");
   S.Diag(FD->getLocation(), diag::err_sycl_mismatch_group_size)
       << /*undefined SYCL_EXTERNAL*/ 1;
diff --git a/clang/test/SemaSYCL/sub-group-size.cpp b/clang/test/SemaSYCL/sub-group-size.cpp
index 414d4b10159e2..112a507f894ba 100644
--- a/clang/test/SemaSYCL/sub-group-size.cpp
+++ b/clang/test/SemaSYCL/sub-group-size.cpp
@@ -142,3 +142,10 @@ void calls_kernel_5() {
 
   });
 }
+
+// Don't diag with the old sub-group-size.
+void calls_kernel_6() {
+  sycl::kernel_single_task<class Kernel6>([]() [[intel::reqd_sub_group_size(10)]] { // #Kernel6
+    NoAttrExternalNotDefined();
+  });
+}

From 51f761733bba86e1463d4f90908b3fd8b11d3665 Mon Sep 17 00:00:00 2001
From: Erich Keane <erich.keane@intel.com>
Date: Wed, 14 Apr 2021 09:34:15 -0700
Subject: [PATCH 22/22] Add comments to each test

---
 clang/test/CodeGenSYCL/sub-group-size.cpp | 3 +++
 clang/test/Driver/sycl-sub-group-size.cpp | 4 ++++
 clang/test/SemaSYCL/sub-group-size.cpp    | 6 ++++++
 3 files changed, 13 insertions(+)

diff --git a/clang/test/CodeGenSYCL/sub-group-size.cpp b/clang/test/CodeGenSYCL/sub-group-size.cpp
index 2f7be0d4c1d20..5df720dbf586c 100644
--- a/clang/test/CodeGenSYCL/sub-group-size.cpp
+++ b/clang/test/CodeGenSYCL/sub-group-size.cpp
@@ -2,6 +2,9 @@
 // RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -fsycl-default-sub-group-size=primary -sycl-std=2020 -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s --check-prefixes=PRIM_DEF,ALL
 // RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -fsycl-default-sub-group-size=10 -sycl-std=2020 -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s --check-prefixes=TEN_DEF,ALL
 
+// Ensure that both forms of the new sub_group_size properly emit their metadata
+// on sycl-kernel and sycl-external functions.
+
 #include "Inputs/sycl.hpp"
 using namespace cl::sycl;
 
diff --git a/clang/test/Driver/sycl-sub-group-size.cpp b/clang/test/Driver/sycl-sub-group-size.cpp
index 151ba82fc873e..a5f85c6581f97 100644
--- a/clang/test/Driver/sycl-sub-group-size.cpp
+++ b/clang/test/Driver/sycl-sub-group-size.cpp
@@ -1,6 +1,10 @@
+// Ensure that by default the -fsycl-default-sub-group-size doesn't get passed
+// to the cc1 invocation.
 // RUN: %clang -### -fsycl %s 2>&1 | FileCheck %s -check-prefix=CHECK-DEFAULT
 // CHECK-DEFAULT-NOT: "fsycl-default-sub-group-size"
 
+// The next two tests make sure that the -fsycl-default-sub-group-size command
+// line arguments get properly passed unaltered to the cc1 invocation.
 // RUN: %clang -### -fsycl -fsycl-default-sub-group-size=primary %s 2>&1 | FileCheck %s -check-prefix=PRIM
 // PRIM: "-fsycl-default-sub-group-size" "primary"
 
diff --git a/clang/test/SemaSYCL/sub-group-size.cpp b/clang/test/SemaSYCL/sub-group-size.cpp
index 112a507f894ba..6578e9e41ee2a 100644
--- a/clang/test/SemaSYCL/sub-group-size.cpp
+++ b/clang/test/SemaSYCL/sub-group-size.cpp
@@ -2,6 +2,12 @@
 // RUN: %clang_cc1 -internal-isystem %S/Inputs -fsycl-is-device -fsycl-default-sub-group-size=primary -sycl-std=2020 -internal-isystem %S/Inputs -fsyntax-only -verify=expected,integer %s
 // RUN: %clang_cc1 -internal-isystem %S/Inputs -fsycl-is-device -fsycl-default-sub-group-size=10 -sycl-std=2020 -internal-isystem %S/Inputs -fsyntax-only -verify=expected,primary %s
 
+// Validate the semantic analysis checks for the interaction betwen the
+// named_sub_group_size and sub_group_size attributes. These are not able to be
+// combined, and require that they only be applied to non-sycl-kernel/
+// non-sycl-device functions if they match the kernel they are being called
+// from.
+
 #include "Inputs/sycl.hpp"
 
 // expected-error@+2 {{'named_sub_group_size' and 'sub_group_size' attributes are not compatible}}