Revert "[OpenMP] Add the ompx_attribute clause for target directives"

This reverts commit ef9ec4b.

The changes broke several bots:
https://lab.llvm.org/buildbot/#/builders/176/builds/3408
https://lab.llvm.org/buildbot/#/builders/198/builds/4028
https://lab.llvm.org/buildbot/#/builders/197/builds/8491
https://lab.llvm.org/buildbot/#/builders/197/builds/8491

Author: AaronBallman

clang/include/clang/AST/OpenMPClause.h

@@ -9172,54 +9172,6 @@ class OMPDoacrossClause final
   }
 };
 
-/// This represents 'ompx_attribute' clause in a directive that might generate
-/// an outlined function. An example is given below.
-///
-/// \code
-/// #pragma omp target [...] ompx_attribute(flatten)
-/// \endcode
-class OMPXAttributeClause
-    : public OMPNoChildClause<llvm::omp::OMPC_ompx_attribute> {
-  friend class OMPClauseReader;
-
-  /// Location of '('.
-  SourceLocation LParenLoc;
-
-  /// The parsed attributes (clause arguments)
-  SmallVector<const Attr *> Attrs;
-
-public:
-  /// Build 'ompx_attribute' clause.
-  ///
-  /// \param Attrs The parsed attributes (clause arguments)
-  /// \param StartLoc Starting location of the clause.
-  /// \param LParenLoc Location of '('.
-  /// \param EndLoc Ending location of the clause.
-  OMPXAttributeClause(ArrayRef<const Attr *> Attrs, SourceLocation StartLoc,
-                      SourceLocation LParenLoc, SourceLocation EndLoc)
-      : OMPNoChildClause(StartLoc, EndLoc), LParenLoc(LParenLoc), Attrs(Attrs) {
-  }
-
-  /// Build an empty clause.
-  OMPXAttributeClause() : OMPNoChildClause() {}
-
-  /// Sets the location of '('.
-  void setLParenLoc(SourceLocation Loc) { LParenLoc = Loc; }
-
-  /// Returns the location of '('.
-  SourceLocation getLParenLoc() const { return LParenLoc; }
-
-  /// Returned the attributes parsed from this clause.
-  ArrayRef<const Attr *> getAttrs() const { return Attrs; }
-
-private:
-  /// Replace the attributes with \p NewAttrs.
-  void setAttrs(ArrayRef<Attr *> NewAttrs) {
-    Attrs.clear();
-    Attrs.append(NewAttrs.begin(), NewAttrs.end());
-  }
-};
-
 } // namespace clang
 
 #endif // LLVM_CLANG_AST_OPENMPCLAUSE_H

clang/include/clang/AST/RecursiveASTVisitor.h

@@ -3884,12 +3884,6 @@ bool RecursiveASTVisitor<Derived>::VisitOMPDoacrossClause(
   return true;
 }
 
-template <typename Derived>
-bool RecursiveASTVisitor<Derived>::VisitOMPXAttributeClause(
-    OMPXAttributeClause *C) {
-  return true;
-}
-
 // FIXME: look at the following tricky-seeming exprs to see if we
 // need to recurse on anything.  These are ones that have methods
 // returning decls or qualtypes or nestednamespecifier -- though I'm

clang/include/clang/Basic/DiagnosticGroups.td

@@ -1308,10 +1308,9 @@ def OpenMPMapping : DiagGroup<"openmp-mapping">;
 def OpenMPTarget : DiagGroup<"openmp-target", [OpenMPMapping]>;
 def OpenMPPre51Compat : DiagGroup<"pre-openmp-51-compat">;
 def OpenMP51Ext : DiagGroup<"openmp-51-extensions">;
-def OpenMPExtensions : DiagGroup<"openmp-extensions">;
 def OpenMP : DiagGroup<"openmp", [
     SourceUsesOpenMP, OpenMPClauses, OpenMPLoopForm, OpenMPTarget,
-    OpenMPMapping, OpenMP51Ext, OpenMPExtensions
+    OpenMPMapping, OpenMP51Ext
   ]>;
 
 // Backend warnings.

clang/include/clang/Basic/DiagnosticParseKinds.td

@@ -1549,9 +1549,6 @@ def warn_omp_more_one_omp_all_memory : Warning<
   InGroup<OpenMPClauses>;
 def warn_omp_depend_in_ordered_deprecated : Warning<"'depend' clause for"
   " 'ordered' is deprecated; use 'doacross' instead">, InGroup<Deprecated>;
-def warn_omp_invalid_attribute_for_ompx_attributes : Warning<"'ompx_attribute' clause only allows "
-  "'amdgpu_flat_work_group_size', 'amdgpu_waves_per_eu', and 'launch_bounds'; "
-  "%0 is ignored">, InGroup<OpenMPExtensions>;
 
 // Pragma loop support.
 def err_pragma_loop_missing_argument : Error<

clang/include/clang/Parse/Parser.h

@@ -3493,13 +3493,6 @@ class Parser : public CodeCompletionHandler {
   //
   OMPClause *ParseOpenMPInteropClause(OpenMPClauseKind Kind, bool ParseOnly);
 
-  /// Parses a ompx_attribute clause
-  ///
-  /// \param ParseOnly true to skip the clause's semantic actions and return
-  /// nullptr.
-  //
-  OMPClause *ParseOpenMPOMPXAttributesClause(bool ParseOnly);
-
 public:
   /// Parses simple expression in parens for single-expression clauses of OpenMP
   /// constructs.

clang/include/clang/Sema/Sema.h

@@ -11034,11 +11034,6 @@ class Sema final {
   bool ConstantFoldAttrArgs(const AttributeCommonInfo &CI,
                             MutableArrayRef<Expr *> Args);
 
-  /// Create an CUDALaunchBoundsAttr attribute.
-  CUDALaunchBoundsAttr *CreateLaunchBoundsAttr(const AttributeCommonInfo &CI,
-                                               Expr *MaxThreads,
-                                               Expr *MinBlocks);
-
   /// AddLaunchBoundsAttr - Adds a launch_bounds attribute to a particular
   /// declaration.
   void AddLaunchBoundsAttr(Decl *D, const AttributeCommonInfo &CI,
@@ -11055,21 +11050,11 @@ class Sema final {
   void AddXConsumedAttr(Decl *D, const AttributeCommonInfo &CI,
                         RetainOwnershipKind K, bool IsTemplateInstantiation);
 
-  /// Create an AMDGPUWavesPerEUAttr attribute.
-  AMDGPUFlatWorkGroupSizeAttr *
-  CreateAMDGPUFlatWorkGroupSizeAttr(const AttributeCommonInfo &CI, Expr *Min,
-                                    Expr *Max);
-
   /// addAMDGPUFlatWorkGroupSizeAttr - Adds an amdgpu_flat_work_group_size
   /// attribute to a particular declaration.
   void addAMDGPUFlatWorkGroupSizeAttr(Decl *D, const AttributeCommonInfo &CI,
                                       Expr *Min, Expr *Max);
 
-  /// Create an AMDGPUWavesPerEUAttr attribute.
-  AMDGPUWavesPerEUAttr *
-  CreateAMDGPUWavesPerEUAttr(const AttributeCommonInfo &CI, Expr *Min,
-                             Expr *Max);
-
   /// addAMDGPUWavePersEUAttr - Adds an amdgpu_waves_per_eu attribute to a
   /// particular declaration.
   void addAMDGPUWavesPerEUAttr(Decl *D, const AttributeCommonInfo &CI,
@@ -12402,12 +12387,6 @@ class Sema final {
                             ArrayRef<Expr *> VarList, SourceLocation StartLoc,
                             SourceLocation LParenLoc, SourceLocation EndLoc);
 
-  /// Called on a well-formed 'ompx_attribute' clause.
-  OMPClause *ActOnOpenMPXAttributeClause(ArrayRef<const Attr *> Attrs,
-                                         SourceLocation StartLoc,
-                                         SourceLocation LParenLoc,
-                                         SourceLocation EndLoc);
-
   /// The kind of conversion being performed.
   enum CheckedConversionKind {
     /// An implicit conversion.

clang/lib/AST/OpenMPClause.cpp

@@ -2534,18 +2534,6 @@ void OMPClausePrinter::VisitOMPDoacrossClause(OMPDoacrossClause *Node) {
   OS << ")";
 }
 
-void OMPClausePrinter::VisitOMPXAttributeClause(OMPXAttributeClause *Node) {
-  OS << "ompx_attribute(";
-  bool IsFirst = true;
-  for (auto &Attr : Node->getAttrs()) {
-    if (!IsFirst)
-      OS << ", ";
-    Attr->printPretty(OS, Policy);
-    IsFirst = false;
-  }
-  OS << ")";
-}
-
 void OMPTraitInfo::getAsVariantMatchInfo(ASTContext &ASTCtx,
                                          VariantMatchInfo &VMI) const {
   for (const OMPTraitSet &Set : Sets) {

clang/lib/AST/StmtProfile.cpp

@@ -928,8 +928,6 @@ void OMPClauseProfiler::VisitOMPXDynCGroupMemClause(
 void OMPClauseProfiler::VisitOMPDoacrossClause(const OMPDoacrossClause *C) {
   VisitOMPClauseList(C);
 }
-void OMPClauseProfiler::VisitOMPXAttributeClause(const OMPXAttributeClause *C) {
-}
 } // namespace
 
 void

clang/lib/CodeGen/CGOpenMPRuntime.cpp

@@ -6115,23 +6115,8 @@ void CGOpenMPRuntime::emitTargetOutlinedFunctionHelper(
                                       DefaultValTeams, DefaultValThreads,
                                       IsOffloadEntry, OutlinedFn, OutlinedFnID);
 
-  if (!OutlinedFn)
-    return;
-
-  CGM.getTargetCodeGenInfo().setTargetAttributes(nullptr, OutlinedFn, CGM);
-
-  for (auto *C : D.getClausesOfKind<OMPXAttributeClause>()) {
-    for (auto *A : C->getAttrs()) {
-      if (auto *Attr = dyn_cast<CUDALaunchBoundsAttr>(A))
-        CGM.handleCUDALaunchBoundsAttr(OutlinedFn, Attr);
-      else if (auto *Attr = dyn_cast<AMDGPUFlatWorkGroupSizeAttr>(A))
-        CGM.handleAMDGPUFlatWorkGroupSizeAttr(OutlinedFn, Attr);
-      else if (auto *Attr = dyn_cast<AMDGPUWavesPerEUAttr>(A))
-        CGM.handleAMDGPUWavesPerEUAttr(OutlinedFn, Attr);
-      else
-        llvm_unreachable("Unexpected attribute kind");
-    }
-  }
+  if (OutlinedFn != nullptr)
+    CGM.getTargetCodeGenInfo().setTargetAttributes(nullptr, OutlinedFn, CGM);
 }
 
 /// Checks if the expression is constant or does not have non-trivial function

clang/lib/CodeGen/CodeGenModule.h

@@ -1588,21 +1588,6 @@ class CodeGenModule : public CodeGenTypeCache {
   /// because we'll lose all important information after each repl.
   void moveLazyEmissionStates(CodeGenModule *NewBuilder);
 
-  /// Emit the IR encoding to attach the CUDA launch bounds attribute to \p F.
-  void handleCUDALaunchBoundsAttr(llvm::Function *F,
-                                  const CUDALaunchBoundsAttr *A);
-
-  /// Emit the IR encoding to attach the AMD GPU flat-work-group-size attribute
-  /// to \p F. Alternatively, the work group size can be taken from a \p
-  /// ReqdWGS.
-  void handleAMDGPUFlatWorkGroupSizeAttr(
-      llvm::Function *F, const AMDGPUFlatWorkGroupSizeAttr *A,
-      const ReqdWorkGroupSizeAttr *ReqdWGS = nullptr);
-
-  /// Emit the IR encoding to attach the AMD GPU waves-per-eu attribute to \p F.
-  void handleAMDGPUWavesPerEUAttr(llvm::Function *F,
-                                  const AMDGPUWavesPerEUAttr *A);
-
 private:
   llvm::Constant *GetOrCreateLLVMFunction(
       StringRef MangledName, llvm::Type *Ty, GlobalDecl D, bool ForVTable,