[clang] Add #pragma clang loop [no]prefetch() (part1)

[Flightor]

Summary

Useage of pragma:
1. #pragma clang loop noprefetch(variable)
   1) `variable`:
      a memory reference(data to be prefetched), must be a declared
      pointer/array variable.

2. #pragma clang loop prefetch(variable, level, distance)
   1) `variable`:
      a memory reference(data to be prefetched), must be a declared
      pointer/array variable.
   2) `level`:
      an optional value to the compiler to specify the type of prefetch.
      '0': data will not be reused
      '1': L1 cache
      '2': L2 cache
      '3': L3 cache
      To use this argument, you must also specify `variable`.
   3) `distance`:
      an option integer argument with a value greater than 0. It indicates
      the number of loop iterations ahead of which a prefetch is issued,
      before the corresponding load or store instruction.
      To use this argument, you must also specify `variable` and `level`.

And add new medatada to support prefetch:

!llvm.loop.prefetch !0
!0 = distinct !{i1 true/false, i32 level, i32 distance}

1st arg: true: prefetch;  false: noprefetch
2nd arg: -1: default;     0-3: cache level
3rd arg: -1: default;     1-INT_MAX: iterations ahead

[Meinersbur]

Can we also see the code that makes use of the new metadata?

clang/include/clang/Parse/LoopHint.h

40

The name and its comment is pretty undescriptive.

clang/lib/CodeGen/CGLoopInfo.cpp

664

[style] LLVM does not use Almost-Always-Auto.

670

[style] LLVM does not use Almost-Always-Auto.

671

[style] If using auto, still use * and &.

Please also consider the LLVM coding style for your other code as well.

892

The metadata must be documented at https://llvm.org/docs/LangRef.html, not just in a hard to find code comment.

The metadata itself seems not be connected to any loop, just to memory access instructions. Since instructions can be moved out of or into a loop, it might get associated with the wrong loop.

clang/lib/CodeGen/CGLoopInfo.h

88

With this many tuple arguments, consider using a struct instead, documenting the meaning of its values.

clang/lib/Parse/ParsePragma.cpp

1459

ParseConstantExpression() can overwrite the Tok field, it must be checked whether it really is a comma before any further parsing.

clang/test/CodeGenCXX/pragma-prefetch-noprefetch.cpp

21

Shouldn't a comma without expression result in an error?

52

167–168

Combining OpenMP with loop hints is unsupported, I don't think it needs to be tested.

[Flightor]

clang/lib/CodeGen/CGLoopInfo.cpp

892

https://reviews.llvm.org/D144378

And some background info: https://discourse.llvm.org/t/rfc-adding-support-pragma-clang-loop-no-prefetch-for-prefetch/68597.

Yeah, it is not bound to the loop at present. You are right. I only deal with this situation in LICM. It is indeed best to add another metadata to bind to the loop id.

[Flightor]

clang/lib/Parse/ParsePragma.cpp

1459

Do you mean that the Tok before line 1458 must be a comma? If so, then the check is made on line 1451.

clang/test/CodeGenCXX/pragma-prefetch-noprefetch.cpp

21

Because noprefetch only accepts one argument (different from prefetch), when the written argument is greater than the specified number of argument, the redundant information is automatically discarded. Referring to the design of other #pragma clang loop *, the alarm information is judged as warning.

167–168

For the message "Combining OpenMP with loop hints is unsupported", I have not found the description of LLVM. Where can I find this message? Thank you.

[congzhe]

clang/lib/CodeGen/CGLoopInfo.cpp

901

is ptr %0 supposed to be ptr %A?

clang/lib/CodeGen/CGLoopInfo.h

38

The comment here seems to be a bit vague? It does not look clear what this value is -- essentially every value either "needs to be prefetched" or "does not need to be prefetched."

Would it be better to change it to something like The value that will be prefetched with "#pragma clang loop prefetch", or refrained from being prefetched with "#pragma clang loop noprefetch"?

88

Would it be better to use PrefetchHint* as elements in the vector instead of using PrefetchHint?

clang/test/CodeGenCXX/pragma-prefetch-noprefetch.cpp

2

Would it be worthwhile adding test cases that have nested loops and with pragma only for the inner loop?

[Meinersbur]

clang/lib/Parse/ParsePragma.cpp

1459

I did not take into account that ParseConstantExpression() uses Tok as the first token. But it means that the comment // ',' is incorrect, it does not parse a comma, it parses the token after the comma and discards the comma that Tok was before. Consider using ConsumeToken or ExpectAndConsume instead.

clang/test/CodeGenCXX/pragma-prefetch-noprefetch.cpp

21

This is probably an oversight, the message warning: extra tokens at end of '#pragma clang loop noprefetch' - ignored is at the very least misleeding:

#pragma clang loop noprefetch(a,) vectorize_enable(4)
for (int i = 0; i < n; ++i)

Is the vectorize_enable taken into account or not?

You also get it not just with additional argument, but anything following after parsing the constant expression:

#pragma clang loop unroll_count(4 4)
for (int i = 0; i < n; ++i)

167–168

It's not a message, it's just nobody has even considered it and hence undocumented. At the very least, the meaning is ambiguous. #pragma omp for may split the loop into two different loops, a "chunk"-loop and an outer loop over chunks and it is unclear to which it should even apply.

I tested it myself some time and Clang would just drop some metadata. I am surprised the !llvm.loop.prefetch is even there.

[Flightor]

Thank for everyone's suggestions. I am redesigning the pragma and clarifying the application scenarios of the pragma.
More discusssion: https://discourse.llvm.org/t/rfc-adding-support-pragma-clang-loop-no-prefetch-for-prefetch/68597/9
A new patch (D146122) to is used to support #pragma clang loop prefetch(disable) first.

clang/docs/LanguageExtensions.rst

.. list-table:: Parameters (unsupported by Clang)
    * - y
      - Enable frame pointers
 
 Extensions for loop hint optimizations
 ======================================
 
 The ``#pragma clang loop`` directive is used to specify hints for optimizing the
 subsequent for, while, do-while, or c++11 range-based for loop. The directive
-provides options for vectorization, interleaving, predication, unrolling and
-distribution. Loop hints can be specified before any loop and will be ignored if
-the optimization is not safe to apply.
+provides options for vectorization, interleaving, predication, unrolling, and
+distribution. And the directive provides two extended options for prefetch and
+noprefetch, which do not support do-while. Loop hints can be specified before
+any loop and will be ignored if the optimization is not safe to apply.
 
 There are loop hints that control transformations (e.g. vectorization, loop
 unrolling) and there are loop hints that set transformation options (e.g.
 ``vectorize_width``, ``unroll_count``).  Pragmas setting transformation options
 imply the transformation is enabled, as if it was enabled via the corresponding
 transformation pragma (e.g. ``vectorize(enable)``). If the transformation is
 disabled  (e.g. ``vectorize(disable)``), that takes precedence over
 transformations option pragmas implying that transformation.
for(...) {
     ...
   }
 
 This predicates (masks) all instructions in the loop, which allows the scalar
 remainder loop (the tail) to be folded into the main vectorized loop. This
 might be more efficient when vector predication is efficiently supported by the
 target platform.
 
+
+Prefetch and NoPrefetch
+-----------------------
+
+Prefetch can effectively improve the performance of programs affected by cache miss.
+
+.. code-block:: c++
+
+  #pragma clang loop prefetch(variable[, level[, distance]])
+  for(...) {
+    ...
+  }
+
+There are three arguments:
+1. ``variable``: a memory reference(data to be prefetched), must be a declared
+``pointer``/``array`` variable.
+2. ``level``: an optional value to the compiler to specify the type of prefetch.
+'0': data will not be reused; '1': L1 cache; '2': L2 cache; '3': L3 cache. To use
+this argument, must also specify ``variable``.
+3. ``distance``: an option integer argument with a value greater than 0. It
+indicates the number of loop iterations ahead of which a prefetch is issued,
+before the corresponding load or store instruction. To use this argument,
+must also specify ``variable`` and ``level``.
+
+NoPrefetch can provide better debugging assistance for performance optimization
+of data prefetching.
+
+.. code-block:: c++
+
+  #pragma clang loop noprefetch(variable)
+  for(...) {
+    ...
+  }
+
+There is only one argument:
+``variable``: a memory reference(data to be prefetched), must be a declared
+``pointer``/``array`` variable.
+
+Prefetch information can be specified for different data by using the
+prefetch/noprefetch pragma multiple times.
+
+.. code-block:: c++
+  #pragma clang loop prefetch(variable1)
+  #pragma clang loop prefetch(variable2)
+  #pragma clang loop noprefetch(variable3)
+  for (...) {
+    ...
+  }
+
 Loop Unrolling
 --------------
 
 Unrolling a loop reduces the loop control overhead and exposes more
 opportunities for ILP. Loops can be fully or partially unrolled. Full unrolling
 eliminates the loop and replaces it with an enumerated sequence of loop
 iterations. Full unrolling is only possible if the loop trip count is known at
 compile time. Partial unrolling replicates the loop body within the loop and

clang/include/clang/Basic/Attr.td

let Spellings = [Pragma<"clang", "loop">, Pragma<"", "unroll">,
                    Pragma<"", "nounroll">, Pragma<"", "unroll_and_jam">,
                    Pragma<"", "nounroll_and_jam">];
 
   /// State of the loop optimization specified by the spelling.
   let Args = [EnumArgument<"Option", "OptionType",
                           ["vectorize", "vectorize_width", "interleave", "interleave_count",
                            "unroll", "unroll_count", "unroll_and_jam", "unroll_and_jam_count",
                            "pipeline", "pipeline_initiation_interval", "distribute",
-                           "vectorize_predicate"],
+                           "vectorize_predicate", "prefetch", "noprefetch"],
                           ["Vectorize", "VectorizeWidth", "Interleave", "InterleaveCount",
                            "Unroll", "UnrollCount", "UnrollAndJam", "UnrollAndJamCount",
                            "PipelineDisabled", "PipelineInitiationInterval", "Distribute",
-                           "VectorizePredicate"]>,
+                           "VectorizePredicate", "Prefetch", "NoPrefetch"]>,
               EnumArgument<"State", "LoopHintState",
                            ["enable", "disable", "numeric", "fixed_width",
-                            "scalable_width", "assume_safety", "full"],
+                            "scalable_width", "assume_safety", "full", "Reference"],
                            ["Enable", "Disable", "Numeric", "FixedWidth",
-                            "ScalableWidth", "AssumeSafety", "Full"]>,
-              ExprArgument<"Value">];
+                            "ScalableWidth", "AssumeSafety", "Full", "Reference"]>,
+              ExprArgument<"DeclRef">, ExprArgument<"Value">, ExprArgument<"Value1">];
 
   let AdditionalMembers = [{
   static const char *getOptionName(int Option) {
     switch(Option) {
     case Vectorize: return "vectorize";
     case VectorizeWidth: return "vectorize_width";
     case Interleave: return "interleave";
     case InterleaveCount: return "interleave_count";
     case Unroll: return "unroll";
     case UnrollCount: return "unroll_count";
     case UnrollAndJam: return "unroll_and_jam";
     case UnrollAndJamCount: return "unroll_and_jam_count";
     case PipelineDisabled: return "pipeline";
     case PipelineInitiationInterval: return "pipeline_initiation_interval";
     case Distribute: return "distribute";
     case VectorizePredicate: return "vectorize_predicate";
+    case Prefetch: return "prefetch";
+    case NoPrefetch: return "noprefetch";
     }
     llvm_unreachable("Unhandled LoopHint option.");
   }
 
   void printPrettyPragma(raw_ostream &OS, const PrintingPolicy &Policy) const;
 
   // Return a string containing the loop hint argument including the
   // enclosing parentheses.

clang/include/clang/Basic/DiagnosticParseKinds.td

def err_omp_requires_out_inout_depend_type : Error<
   "reserved locator 'omp_all_memory' requires 'out' or 'inout' "
   "dependency types">;
 def warn_omp_more_one_omp_all_memory : Warning<
   "reserved locator 'omp_all_memory' cannot be specified more than once">,
   InGroup<OpenMPClauses>;
 
 // Pragma loop support.
 def err_pragma_loop_missing_argument : Error<
-  "missing argument; expected %select{an integer value|"
-  "'enable'%select{|, 'full'}1%select{|, 'assume_safety'}2 or 'disable'}0">;
+  "missing argument; expected %select{an %select{integer value|declared variable}1|"
+  "'enable'%select{|, 'full'}2%select{|, 'assume_safety'}3 or 'disable'}0">;
 def err_pragma_loop_invalid_option : Error<
   "%select{invalid|missing}0 option%select{ %1|}0; expected vectorize, "
   "vectorize_width, interleave, interleave_count, unroll, unroll_count, "
-  "pipeline, pipeline_initiation_interval, vectorize_predicate, or distribute">;
+  "pipeline, pipeline_initiation_interval, vectorize_predicate, distribute, "
+  "prefetch, or noprefetch">;
 def err_pragma_loop_invalid_vectorize_option : Error<
   "vectorize_width loop hint malformed; use vectorize_width(X, fixed) or "
   "vectorize_width(X, scalable) where X is an integer, or vectorize_width('fixed' or 'scalable')">;
+def err_pragma_prefetch_invalid_arg : Error<
+  "invalid argument; expected an declared variable">;
+def err_pragma_prefetch_punc : Error<"expected ')' or ','%select{| after '%1'}0">;
 def note_pragma_loop_invalid_vectorize_option : Note<
   "vectorize_width loop hint malformed; use vectorize_width(X, fixed) or "
   "vectorize_width(X, scalable) where X is an integer, or vectorize_width('fixed' or 'scalable')">;
 
 def err_pragma_fp_invalid_option : Error<
   "%select{invalid|missing}0 option%select{ %1|}0; expected 'contract', 'reassociate' or 'exceptions'">;
 def err_pragma_fp_invalid_argument : Error<
   "unexpected argument '%0' to '#pragma clang fp %1'; expected "

clang/include/clang/Basic/DiagnosticSemaKinds.td

 def err_pragma_loop_invalid_argument_type : Error<
   "invalid argument of type %0; expected an integer type">;
 def err_pragma_loop_invalid_argument_value : Error<
   "%select{invalid value '%0'; must be positive|value '%0' is too large}1">;
 def err_pragma_loop_compatibility : Error<
   "%select{incompatible|duplicate}0 directives '%1' and '%2'">;
 def err_pragma_loop_precedes_nonloop : Error<
   "expected a for, while, or do-while loop to follow '%0'">;
+def err_pragma_loop_invalid_numeric : Error<
+  "invalid value '%0'; expected integer constant between 0 and 3">;
 
 def err_pragma_attribute_matcher_subrule_contradicts_rule : Error<
   "redundant attribute subject matcher sub-rule '%0'; '%1' already matches "
   "those declarations">;
 def err_pragma_attribute_matcher_negated_subrule_contradicts_subrule : Error<
   "negated attribute subject matcher sub-rule '%0' contradicts sub-rule '%1'">;
 def err_pragma_attribute_invalid_matchers : Error<
   "attribute %0 can't be applied to %1">;

clang/include/clang/Parse/LoopHint.h

struct LoopHint {
   IdentifierLoc *PragmaNameLoc;
   // Name of the loop hint.  Examples: "unroll", "vectorize".  In the
   // "#pragma unroll" and "#pragma nounroll" cases, this is identical to
   // PragmaNameLoc.
   IdentifierLoc *OptionLoc;
   // Identifier for the hint state argument.  If null, then the state is
   // default value such as for "#pragma unroll".
   IdentifierLoc *StateLoc;
+  // Declaration reference expression for the hint argument if it exits,
+  // null otherwise.
+  Expr *DeclRef;
   // Expression for the hint argument if it exists, null otherwise.
   Expr *ValueExpr;
+  // Expression for the hint argument if it exists, null otherwise.
+  Expr *ValueExpr1;
 
   LoopHint()
       : PragmaNameLoc(nullptr), OptionLoc(nullptr), StateLoc(nullptr),
-        ValueExpr(nullptr) {}
+        DeclRef(nullptr), ValueExpr(nullptr), ValueExpr1(nullptr) {}
 };
 
 } // end namespace clang
 
 #endif // LLVM_CLANG_PARSE_LOOPHINT_H

clang/include/clang/Sema/Sema.h

ExprResult BuildSYCLUniqueStableNameExpr(SourceLocation OpLoc,
                                            SourceLocation LParen,
                                            SourceLocation RParen,
                                            TypeSourceInfo *TSI);
   ExprResult ActOnSYCLUniqueStableNameExpr(SourceLocation OpLoc,
                                            SourceLocation LParen,
                                            SourceLocation RParen,
                                            ParsedType ParsedTy);
 
-  bool CheckLoopHintExpr(Expr *E, SourceLocation Loc);
+  bool CheckLoopHintExpr(Expr *E, SourceLocation Loc, bool IsPrefetch = false);
 
   ExprResult ActOnNumericConstant(const Token &Tok, Scope *UDLScope = nullptr);
   ExprResult ActOnCharacterConstant(const Token &Tok,
                                     Scope *UDLScope = nullptr);
   ExprResult ActOnParenExpr(SourceLocation L, SourceLocation R, Expr *E);
   ExprResult ActOnParenListExpr(SourceLocation L,
                                 SourceLocation R,
                                 MultiExprArg Val);

clang/lib/AST/AttrImpl.cpp

 // Return a string containing the loop hint argument including the
 // enclosing parentheses.
 std::string LoopHintAttr::getValueString(const PrintingPolicy &Policy) const {
   std::string ValueName;
   llvm::raw_string_ostream OS(ValueName);
   OS << "(";
   if (state == Numeric)
     value->printPretty(OS, nullptr, Policy);
-  else if (state == FixedWidth || state == ScalableWidth) {
+  else if (state == Reference) {
+    declRef->printPretty(OS, nullptr, Policy);
+    if (value) {
+      OS << ", ";
+      value->printPretty(OS, nullptr, Policy);
+      if (value1) {
+        OS << ", ";
+        value1->printPretty(OS, nullptr, Policy);
+      }
+    }
+  } else if (state == FixedWidth || state == ScalableWidth) {
     if (value) {
       value->printPretty(OS, nullptr, Policy);
       if (state == ScalableWidth)
         OS << ", scalable";
     } else if (state == ScalableWidth)
       OS << "scalable";
     else
       OS << "fixed";

clang/lib/CodeGen/CGLoopInfo.h

 //===---- CGLoopInfo.h - LLVM CodeGen for loop metadata -*- C++ -*---------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
 // This is the internal state used for llvm translation for loop statement
 // metadata.
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef LLVM_CLANG_LIB_CODEGEN_CGLOOPINFO_H
 #define LLVM_CLANG_LIB_CODEGEN_CGLOOPINFO_H
 
+#include "Address.h"
 #include "llvm/ADT/ArrayRef.h"
 #include "llvm/ADT/SmallVector.h"
 #include "llvm/IR/DebugLoc.h"
 #include "llvm/IR/Value.h"
 #include "llvm/Support/Compiler.h"
 
 namespace llvm {
 class BasicBlock;
 class Instruction;
 class MDNode;
 } // end namespace llvm
 
 namespace clang {
 class Attr;
 class ASTContext;
 class CodeGenOptions;
+class Decl;
 namespace CodeGen {
 
 /// Attributes that may be specified on loops.
 struct LoopAttributes {
   explicit LoopAttributes(bool IsParallel = false);
   void clear();
 
   /// Generate llvm.loop.parallel metadata for loads and stores.
   bool IsParallel;
 
   /// State of loop vectorization or unrolling.
   enum LVEnableState { Unspecified, Enable, Disable, Full };
struct LoopAttributes {
   /// Value for llvm.loop.pipeline.disable metadata.
   bool PipelineDisabled;
 
   /// Value for llvm.loop.pipeline.iicount metadata.
   unsigned PipelineInitiationInterval;
 
   /// Value for whether the loop is required to make progress.
   bool MustProgress;
+
+  /// Value for llvm.loop.prefetch metadata
+  llvm::SmallVector<std::tuple<llvm::Value *, int, int, bool>, 4> PrefetchInfo;
 };
 
 /// Information used when generating a structured loop.
 class LoopInfo {
 public:
   /// Construct a new LoopInfo for the loop with entry Header.
   LoopInfo(llvm::BasicBlock *Header, const LoopAttributes &Attrs,
            const llvm::DebugLoc &StartLoc, const llvm::DebugLoc &EndLoc,
            LoopInfo *Parent);
 
   /// Get the loop id metadata for this loop.
   llvm::MDNode *getLoopID() const { return TempLoopID.get(); }
 
   /// Get the header block of this loop.
   llvm::BasicBlock *getHeader() const { return Header; }
 
   /// Get the set of attributes active for this loop.
   const LoopAttributes &getAttributes() const { return Attrs; }
 
   /// Return this loop's access group or nullptr if it does not have one.
   llvm::MDNode *getAccessGroup() const { return AccGroup; }
 
+  /// Get the MDNode of prefetch variable for this loop.
+  llvm::MDNode *getPrefetchMDNode(llvm::Value *Ad) {
+    auto I = PrefetchMDs.find(Ad);
+    return I != PrefetchMDs.end() ? I->second : nullptr;
+  };
+
   /// Create the loop's metadata. Must be called after its nested loops have
   /// been processed.
   void finish();
 
 private:
   /// Loop ID metadata.
   llvm::TempMDTuple TempLoopID;
   /// Header block of this loop.
   llvm::BasicBlock *Header;
   /// The attributes for this loop.
   LoopAttributes Attrs;
   /// The access group for memory accesses parallel to this loop.
   llvm::MDNode *AccGroup = nullptr;
   /// Start location of this loop.
   llvm::DebugLoc StartLoc;
   /// End location of this loop.
   llvm::DebugLoc EndLoc;
   /// The next outer loop, or nullptr if this is the outermost loop.
   LoopInfo *Parent;
   /// If this loop has unroll-and-jam metadata, this can be set by the inner
   /// loop's LoopInfo to set the llvm.loop.unroll_and_jam.followup_inner
   /// metadata.
   llvm::MDNode *UnrollAndJamInnerFollowup = nullptr;
+  /// Each prefetch variable has unique metadata.
+  llvm::DenseMap<llvm::Value *, llvm::MDNode *> PrefetchMDs;
 
   /// Create a LoopID without any transformations.
   llvm::MDNode *
   createLoopPropertiesMetadata(llvm::ArrayRef<llvm::Metadata *> LoopProperties);
 
   /// Create a LoopID for transformations.
   ///
   /// The methods call each other in case multiple transformations are applied
public:
   /// applied to the loop and then cleared.
   void push(llvm::BasicBlock *Header, const llvm::DebugLoc &StartLoc,
             const llvm::DebugLoc &EndLoc);
 
   /// Begin a new structured loop. Stage attributes from the Attrs list.
   /// The staged attributes are applied to the loop and then cleared.
   void push(llvm::BasicBlock *Header, clang::ASTContext &Ctx,
             const clang::CodeGenOptions &CGOpts,
-            llvm::ArrayRef<const Attr *> Attrs, const llvm::DebugLoc &StartLoc,
-            const llvm::DebugLoc &EndLoc, bool MustProgress = false);
+            llvm::ArrayRef<const Attr *> Attrs,
+            llvm::DenseMap<const Decl *, Address> &LocalDeclMap,
+            const llvm::DebugLoc &StartLoc, const llvm::DebugLoc &EndLoc,
+            bool MustProgress = false);
 
   /// End the current loop.
   void pop();
 
   /// Return the top loop id metadata.
   llvm::MDNode *getCurLoopID() const { return getInfo().getLoopID(); }
 
   /// Return true if the top loop is parallel.
public:
   /// Set the pipeline initiation interval.
   void setPipelineInitiationInterval(unsigned C) {
     StagedAttrs.PipelineInitiationInterval = C;
   }
 
   /// Set no progress for the next loop pushed.
   void setMustProgress(bool P) { StagedAttrs.MustProgress = P; }
 
+  /// Set prefetch value, cache level, iteration ahead for the next loop pushed,
+  /// or set noprefetch value.
+  void setPrefetchInfo(std::tuple<llvm::Value *, int, int, bool> S) {
+    StagedAttrs.PrefetchInfo.emplace_back(S);
+  }
+
 private:
   /// Returns true if there is LoopInfo on the stack.
   bool hasInfo() const { return !Active.empty(); }
   /// Return the LoopInfo for the current loop. HasInfo should be called
   /// first to ensure LoopInfo is present.
   const LoopInfo &getInfo() const { return *Active.back(); }
   /// The set of attributes that will be applied to the next pushed loop.
   LoopAttributes StagedAttrs;
   /// Stack of active loops.
   llvm::SmallVector<std::unique_ptr<LoopInfo>, 4> Active;
 };
 
 } // end namespace CodeGen
 } // end namespace clang
 
 #endif

clang/lib/CodeGen/CGLoopInfo.cpp

void LoopAttributes::clear() {
   VectorizeEnable = LoopAttributes::Unspecified;
   UnrollEnable = LoopAttributes::Unspecified;
   UnrollAndJamEnable = LoopAttributes::Unspecified;
   VectorizePredicateEnable = LoopAttributes::Unspecified;
   DistributeEnable = LoopAttributes::Unspecified;
   PipelineDisabled = false;
   PipelineInitiationInterval = 0;
   MustProgress = false;
+  PrefetchInfo.clear();
 }
 
 LoopInfo::LoopInfo(BasicBlock *Header, const LoopAttributes &Attrs,
                    const llvm::DebugLoc &StartLoc, const llvm::DebugLoc &EndLoc,
                    LoopInfo *Parent)
     : Header(Header), Attrs(Attrs), StartLoc(StartLoc), EndLoc(EndLoc),
       Parent(Parent) {
 
   if (Attrs.IsParallel) {
     // Create an access group for this loop.
     LLVMContext &Ctx = Header->getContext();
     AccGroup = MDNode::getDistinct(Ctx, {});
   }
 
+  if (!Attrs.PrefetchInfo.empty()) {
+    LLVMContext &Ctx = Header->getContext();
+    for (auto S : Attrs.PrefetchInfo) {
+      Value *Val = std::get<0>(S);
+      int CacheLevel = std::get<1>(S);
+      int IterAhead = std::get<2>(S);
+      bool EnablePrefetch = std::get<3>(S);
+
+      Metadata *Vals[] = {ConstantAsMetadata::get(ConstantInt::get(
+                              llvm::Type::getInt1Ty(Ctx), EnablePrefetch)),
+                          ConstantAsMetadata::get(ConstantInt::get(
+                              llvm::Type::getInt32Ty(Ctx), CacheLevel)),
+                          ConstantAsMetadata::get(ConstantInt::get(
+                              llvm::Type::getInt32Ty(Ctx), IterAhead))};
+      MDNode *MD = MDNode::getDistinct(Ctx, Vals);
+      PrefetchMDs[Val] = MD;
+    }
+  }
+
   if (!Attrs.IsParallel && Attrs.VectorizeWidth == 0 &&
       Attrs.VectorizeScalable == LoopAttributes::Unspecified &&
       Attrs.InterleaveCount == 0 && Attrs.UnrollCount == 0 &&
       Attrs.UnrollAndJamCount == 0 && !Attrs.PipelineDisabled &&
       Attrs.PipelineInitiationInterval == 0 &&
       Attrs.VectorizePredicateEnable == LoopAttributes::Unspecified &&
       Attrs.VectorizeEnable == LoopAttributes::Unspecified &&
       Attrs.UnrollEnable == LoopAttributes::Unspecified &&
       Attrs.UnrollAndJamEnable == LoopAttributes::Unspecified &&
       Attrs.DistributeEnable == LoopAttributes::Unspecified && !StartLoc &&
-      !EndLoc && !Attrs.MustProgress)
+      !EndLoc && !Attrs.MustProgress && Attrs.PrefetchInfo.empty())
     return;
 
   TempLoopID = MDNode::getTemporary(Header->getContext(), std::nullopt);
 }
 
 void LoopInfo::finish() {
   // We did not annotate the loop body instructions because there are no
   // attributes for this loop.
Active.emplace_back(
                    Active.empty() ? nullptr : Active.back().get()));
   // Clear the attributes so nested loops do not inherit them.
   StagedAttrs.clear();
 }
 
 void LoopInfoStack::push(BasicBlock *Header, clang::ASTContext &Ctx,
                          const clang::CodeGenOptions &CGOpts,
                          ArrayRef<const clang::Attr *> Attrs,
+                         llvm::DenseMap<const Decl *, Address> &LocalDeclMap,
                          const llvm::DebugLoc &StartLoc,
                          const llvm::DebugLoc &EndLoc, bool MustProgress) {
   // Identify loop hint attributes from Attrs.
   for (const auto *Attr : Attrs) {
     const LoopHintAttr *LH = dyn_cast<LoopHintAttr>(Attr);
     const OpenCLUnrollHintAttr *OpenCLHint =
         dyn_cast<OpenCLUnrollHintAttr>(Attr);
 
     // Skip non loop hint attributes
     if (!LH && !OpenCLHint) {
       continue;
     }
 
     LoopHintAttr::OptionType Option = LoopHintAttr::Unroll;
     LoopHintAttr::LoopHintState State = LoopHintAttr::Disable;
-    unsigned ValueInt = 1;
+    int ValueInt = 1;
+    int ValueInt1 = -1;
+    Value *DeclRef = nullptr;
     // Translate opencl_unroll_hint attribute argument to
     // equivalent LoopHintAttr enums.
     // OpenCL v2.0 s6.11.5:
     // 0 - enable unroll (no argument).
     // 1 - disable unroll.
     // other positive integer n - unroll by n.
     if (OpenCLHint) {
       ValueInt = OpenCLHint->getUnrollHint();
       if (ValueInt == 0) {
         State = LoopHintAttr::Enable;
       } else if (ValueInt != 1) {
         Option = LoopHintAttr::UnrollCount;
         State = LoopHintAttr::Numeric;
       }
     } else if (LH) {
+      Option = LH->getOption();
+      if (Option == LoopHintAttr::Prefetch ||
+          Option == LoopHintAttr::NoPrefetch)
+        ValueInt = -1;
+
       auto *ValueExpr = LH->getValue();
       if (ValueExpr) {
         llvm::APSInt ValueAPS = ValueExpr->EvaluateKnownConstInt(Ctx);
         ValueInt = ValueAPS.getSExtValue();
       }
 
-      Option = LH->getOption();
+      auto *ValueExpr1 = LH->getValue1();
+      if (ValueExpr1) {
+        llvm::APSInt ValueAPS = ValueExpr1->EvaluateKnownConstInt(Ctx);
+        ValueInt1 = ValueAPS.getSExtValue();
+      }
+
+      if (auto *Ref = LH->getDeclRef()) {
+        auto DRE = dyn_cast<DeclRefExpr>(Ref);
+        assert(DRE && "Attribute must hava a valid declaration reference.");
+        if (auto VD = dyn_cast<VarDecl>(DRE->getDecl())) {
+          auto I = LocalDeclMap.find(VD);
+          if (I != LocalDeclMap.end())
+            DeclRef = I->second.getPointer();
+        }
+      }
+
       State = LH->getState();
     }
     switch (State) {
     case LoopHintAttr::Disable:
       switch (Option) {
       case LoopHintAttr::Vectorize:
         // Disable vectorization by specifying a width of 1.
         setVectorizeWidth(1);
case LoopHintAttr::Disable:
       case LoopHintAttr::PipelineDisabled:
         setPipelineDisabled(true);
         break;
       case LoopHintAttr::UnrollCount:
       case LoopHintAttr::UnrollAndJamCount:
       case LoopHintAttr::VectorizeWidth:
       case LoopHintAttr::InterleaveCount:
       case LoopHintAttr::PipelineInitiationInterval:
+      case LoopHintAttr::NoPrefetch:
+      case LoopHintAttr::Prefetch:
         llvm_unreachable("Options cannot be disabled.");
         break;
       }
       break;
     case LoopHintAttr::Enable:
       switch (Option) {
       case LoopHintAttr::Vectorize:
       case LoopHintAttr::Interleave:
case LoopHintAttr::Enable:
         setDistributeState(true);
         break;
       case LoopHintAttr::UnrollCount:
       case LoopHintAttr::UnrollAndJamCount:
       case LoopHintAttr::VectorizeWidth:
       case LoopHintAttr::InterleaveCount:
       case LoopHintAttr::PipelineDisabled:
       case LoopHintAttr::PipelineInitiationInterval:
+      case LoopHintAttr::NoPrefetch:
+      case LoopHintAttr::Prefetch:
         llvm_unreachable("Options cannot enabled.");
         break;
       }
       break;
     case LoopHintAttr::AssumeSafety:
       switch (Option) {
       case LoopHintAttr::Vectorize:
       case LoopHintAttr::Interleave:
         // Apply "llvm.mem.parallel_loop_access" metadata to load/stores.
         setParallel(true);
         setVectorizeEnable(true);
         break;
       case LoopHintAttr::Unroll:
       case LoopHintAttr::UnrollAndJam:
       case LoopHintAttr::VectorizePredicate:
       case LoopHintAttr::UnrollCount:
       case LoopHintAttr::UnrollAndJamCount:
       case LoopHintAttr::VectorizeWidth:
       case LoopHintAttr::InterleaveCount:
       case LoopHintAttr::Distribute:
       case LoopHintAttr::PipelineDisabled:
       case LoopHintAttr::PipelineInitiationInterval:
+      case LoopHintAttr::NoPrefetch:
+      case LoopHintAttr::Prefetch:
         llvm_unreachable("Options cannot be used to assume mem safety.");
         break;
       }
       break;
     case LoopHintAttr::Full:
       switch (Option) {
       case LoopHintAttr::Unroll:
         setUnrollState(LoopAttributes::Full);
         break;
       case LoopHintAttr::UnrollAndJam:
         setUnrollAndJamState(LoopAttributes::Full);
         break;
       case LoopHintAttr::Vectorize:
       case LoopHintAttr::Interleave:
       case LoopHintAttr::UnrollCount:
       case LoopHintAttr::UnrollAndJamCount:
       case LoopHintAttr::VectorizeWidth:
       case LoopHintAttr::InterleaveCount:
       case LoopHintAttr::Distribute:
       case LoopHintAttr::PipelineDisabled:
       case LoopHintAttr::PipelineInitiationInterval:
       case LoopHintAttr::VectorizePredicate:
+      case LoopHintAttr::NoPrefetch:
+      case LoopHintAttr::Prefetch:
         llvm_unreachable("Options cannot be used with 'full' hint.");
         break;
       }
       break;
     case LoopHintAttr::FixedWidth:
     case LoopHintAttr::ScalableWidth:
       switch (Option) {
       case LoopHintAttr::VectorizeWidth:
case LoopHintAttr::Numeric:
       case LoopHintAttr::Unroll:
       case LoopHintAttr::UnrollAndJam:
       case LoopHintAttr::VectorizePredicate:
       case LoopHintAttr::Vectorize:
       case LoopHintAttr::VectorizeWidth:
       case LoopHintAttr::Interleave:
       case LoopHintAttr::Distribute:
       case LoopHintAttr::PipelineDisabled:
+      case LoopHintAttr::NoPrefetch:
+      case LoopHintAttr::Prefetch:
         llvm_unreachable("Options cannot be assigned a value.");
         break;
       }
       break;
+    case LoopHintAttr::Reference:
+      switch (Option) {
+      case LoopHintAttr::Prefetch:
+        setPrefetchInfo(std::make_tuple(DeclRef, ValueInt, ValueInt1, true));
+        break;
+      case LoopHintAttr::NoPrefetch:
+        setPrefetchInfo(std::make_tuple(DeclRef, ValueInt, ValueInt1, false));
+        break;
+      default:
+        llvm_unreachable("Option cannot be assigned a variable");
+        break;
+      }
+      break;
     }
   }
 
   setMustProgress(MustProgress);
 
   if (CGOpts.OptimizationLevel > 0)
     // Disable unrolling for the loop, if unrolling is disabled (via
     // -fno-unroll-loops) and no pragmas override the decision.
for (const auto &AL : Active) {
         AccessGroups.push_back(Group);
     }
     MDNode *UnionMD = nullptr;
     if (AccessGroups.size() == 1)
       UnionMD = cast<MDNode>(AccessGroups[0]);
     else if (AccessGroups.size() >= 2)
       UnionMD = MDNode::get(I->getContext(), AccessGroups);
     I->setMetadata("llvm.access.group", UnionMD);
+
+    // Add "llvm.loop.prefetch" metadata for load/store of variables that need
+    // to be (no)prefetched.
+    // 1. Assume the prefetch variable is a pointer \p A.
+    //   %A = alloca ptr
+    //   %0 = load ptr, ptr %A
+    //   %arrayidx = getelementptr inbounds i32, ptr %0, i64 %idxprom
+    //   %1 = load i32, ptr %arrayidx
+    // 2. Assume the prefetch variable is an array \p A.
+    //   %A = alloca [n * i32]
+    //   %arrayidx = getelementptr inbounds i32, ptr %0, i64 %idxprom
+    //   %0 = load i32, ptr %arrayidx
+    auto LI = dyn_cast<LoadInst>(I);
+    auto SI = dyn_cast<StoreInst>(I);
+    if (!Active.empty() && (LI || SI)) {
+      auto PrefetchOp = LI ? LI->getPointerOperand() : SI->getPointerOperand();
+      if (auto GEP = dyn_cast<GetElementPtrInst>(PrefetchOp)) {
+        auto AddrLI = dyn_cast<LoadInst>(GEP->getPointerOperand());
+        auto PrefetchVal =
+            AddrLI ? AddrLI->getPointerOperand() : GEP->getPointerOperand();
+        if (PrefetchVal) {
+          MDNode *PrefetchMD = nullptr;
+          // For nested loops, use the prefetch pragma of the outer loop
+          // closest to the current loop.
+          for (const auto &AL : Active) {
+            auto MD = AL->getPrefetchMDNode(PrefetchVal);
+            PrefetchMD = MD ? MD : PrefetchMD;
+          }
+          if (PrefetchMD)
+            I->setMetadata("llvm.loop.prefetch", PrefetchMD);
+        }
+      }
+    }
   }
 
   if (!hasInfo())
     return;
 
   const LoopInfo &L = getInfo();
   if (!L.getLoopID())
     return;

clang/lib/CodeGen/CGStmt.cpp

void CodeGenFunction::EmitWhileStmt(const WhileStmt &S,
 
   // while(1) is common, avoid extra exit blocks.  Be sure
   // to correctly handle break/continue though.
   llvm::ConstantInt *C = dyn_cast<llvm::ConstantInt>(BoolCondVal);
   bool CondIsConstInt = C != nullptr;
   bool EmitBoolCondBranch = !CondIsConstInt || !C->isOne();
   const SourceRange &R = S.getSourceRange();
   LoopStack.push(LoopHeader.getBlock(), CGM.getContext(), CGM.getCodeGenOpts(),
-                 WhileAttrs, SourceLocToDebugLoc(R.getBegin()),
+                 WhileAttrs, LocalDeclMap, SourceLocToDebugLoc(R.getBegin()),
                  SourceLocToDebugLoc(R.getEnd()),
                  checkIfLoopMustProgress(CondIsConstInt));
 
   // As long as the condition is true, go to the loop body.
   llvm::BasicBlock *LoopBody = createBasicBlock("while.body");
   if (EmitBoolCondBranch) {
     llvm::BasicBlock *ExitBlock = LoopExit.getBlock();
     if (ConditionScope.requiresCleanups())
void CodeGenFunction::EmitDoStmt(const DoStmt &S,
   // "do {} while (0)" is common in macros, avoid extra blocks.  Be sure
   // to correctly handle break/continue though.
   llvm::ConstantInt *C = dyn_cast<llvm::ConstantInt>(BoolCondVal);
   bool CondIsConstInt = C;
   bool EmitBoolCondBranch = !C || !C->isZero();
 
   const SourceRange &R = S.getSourceRange();
   LoopStack.push(LoopBody, CGM.getContext(), CGM.getCodeGenOpts(), DoAttrs,
-                 SourceLocToDebugLoc(R.getBegin()),
+                 LocalDeclMap, SourceLocToDebugLoc(R.getBegin()),
                  SourceLocToDebugLoc(R.getEnd()),
                  checkIfLoopMustProgress(CondIsConstInt));
 
   // As long as the condition is true, iterate the loop.
   if (EmitBoolCondBranch) {
     uint64_t BackedgeCount = getProfileCount(S.getBody()) - ParentCount;
     Builder.CreateCondBr(
         BoolCondVal, LoopBody, LoopExit.getBlock(),
void CodeGenFunction::EmitForStmt(const ForStmt &S,
   EmitBlock(CondBlock);
 
   Expr::EvalResult Result;
   bool CondIsConstInt =
       !S.getCond() || S.getCond()->EvaluateAsInt(Result, getContext());
 
   const SourceRange &R = S.getSourceRange();
   LoopStack.push(CondBlock, CGM.getContext(), CGM.getCodeGenOpts(), ForAttrs,
-                 SourceLocToDebugLoc(R.getBegin()),
+                 LocalDeclMap, SourceLocToDebugLoc(R.getBegin()),
                  SourceLocToDebugLoc(R.getEnd()),
                  checkIfLoopMustProgress(CondIsConstInt));
 
   // Create a cleanup scope for the condition variable cleanups.
   LexicalScope ConditionScope(*this, S.getSourceRange());
 
   // If the for loop doesn't have an increment we can just use the condition as
   // the continue block. Otherwise, if there is no condition variable, we can
CodeGenFunction::EmitCXXForRangeStmt(const CXXForRangeStmt &S,
   // Start the loop with a block that tests the condition.
   // If there's an increment, the continue scope will be overwritten
   // later.
   llvm::BasicBlock *CondBlock = createBasicBlock("for.cond");
   EmitBlock(CondBlock);
 
   const SourceRange &R = S.getSourceRange();
   LoopStack.push(CondBlock, CGM.getContext(), CGM.getCodeGenOpts(), ForAttrs,
-                 SourceLocToDebugLoc(R.getBegin()),
+                 LocalDeclMap, SourceLocToDebugLoc(R.getBegin()),
                  SourceLocToDebugLoc(R.getEnd()));
 
   // If there are any cleanups between here and the loop-exit scope,
   // create a block to stage a loop exit along.
   llvm::BasicBlock *ExitBlock = LoopExit.getBlock();
   if (ForScope.requiresCleanups())
     ExitBlock = createBasicBlock("for.cond.cleanup");
 

clang/lib/CodeGen/CGStmtOpenMP.cpp

void CodeGenFunction::EmitOMPInnerLoop(
   // If attributes are attached, push to the basic block with them.
   const auto &OMPED = cast<OMPExecutableDirective>(S);
   const CapturedStmt *ICS = OMPED.getInnermostCapturedStmt();
   const Stmt *SS = ICS->getCapturedStmt();
   const AttributedStmt *AS = dyn_cast_or_null<AttributedStmt>(SS);
   OMPLoopNestStack.clear();
   if (AS)
     LoopStack.push(CondBlock, CGM.getContext(), CGM.getCodeGenOpts(),
-                   AS->getAttrs(), SourceLocToDebugLoc(R.getBegin()),
+                   AS->getAttrs(), LocalDeclMap,
+                   SourceLocToDebugLoc(R.getBegin()),
                    SourceLocToDebugLoc(R.getEnd()));
   else
     LoopStack.push(CondBlock, SourceLocToDebugLoc(R.getBegin()),
                    SourceLocToDebugLoc(R.getEnd()));
 
   // If there are any cleanups between here and the loop-exit scope,
   // create a block to stage a loop exit along.
   llvm::BasicBlock *ExitBlock = LoopExit.getBlock();

clang/lib/Parse/ParsePragma.cpp

bool Parser::HandlePragmaLoopHint(LoopHint &Hint) {
   assert(!Toks.empty() &&
          "PragmaLoopHintInfo::Toks must contain at least one token.");
 
   // If no option is specified the argument is assumed to be a constant expr.
   bool OptionUnroll = false;
   bool OptionUnrollAndJam = false;
   bool OptionDistribute = false;
   bool OptionPipelineDisabled = false;
+  bool OptionNoPrefetch = false;
+  bool OptionPrefetch = false;
   bool StateOption = false;
   if (OptionInfo) { // Pragma Unroll does not specify an option.
     OptionUnroll = OptionInfo->isStr("unroll");
     OptionUnrollAndJam = OptionInfo->isStr("unroll_and_jam");
     OptionDistribute = OptionInfo->isStr("distribute");
     OptionPipelineDisabled = OptionInfo->isStr("pipeline");
+    OptionPrefetch = OptionInfo->isStr("prefetch");
+    OptionNoPrefetch = OptionInfo->isStr("noprefetch");
     StateOption = llvm::StringSwitch<bool>(OptionInfo->getName())
                       .Case("vectorize", true)
                       .Case("interleave", true)
                       .Case("vectorize_predicate", true)
                       .Default(false) ||
                   OptionUnroll || OptionUnrollAndJam || OptionDistribute ||
                   OptionPipelineDisabled;
   }
 
   bool AssumeSafetyArg = !OptionUnroll && !OptionUnrollAndJam &&
                          !OptionDistribute && !OptionPipelineDisabled;
   // Verify loop hint has an argument.
   if (Toks[0].is(tok::eof)) {
     ConsumeAnnotationToken();
     Diag(Toks[0].getLocation(), diag::err_pragma_loop_missing_argument)
         << /*StateArgument=*/StateOption
+        << /*VariableArg=*/(OptionPrefetch || OptionNoPrefetch)
         << /*FullKeyword=*/(OptionUnroll || OptionUnrollAndJam)
         << /*AssumeSafetyKeyword=*/AssumeSafetyArg;
     return false;
   }
 
   // Validate the argument.
   if (StateOption) {
     ConsumeAnnotationToken();
if (!Valid) {
             << /*AssumeSafetyKeyword=*/AssumeSafetyArg;
       }
       return false;
     }
     if (Toks.size() > 2)
       Diag(Tok.getLocation(), diag::warn_pragma_extra_tokens_at_eol)
           << PragmaLoopHintString(Info->PragmaName, Info->Option);
     Hint.StateLoc = IdentifierLoc::create(Actions.Context, StateLoc, StateInfo);
+  } else if (OptionNoPrefetch || OptionPrefetch) {
+    UnqualifiedId Name;
+    CXXScopeSpec ScopeSpec;
+    SourceLocation TemplateKWLoc;
+
+    PP.EnterTokenStream(Toks, /*DisableMacroExpansion=*/false,
+                        /*IsReinject=*/false);
+    ConsumeAnnotationToken();
+
+    IdentifierInfo *StateInfo = Tok.getIdentifierInfo();
+    SourceLocation StateLoc = ConsumeToken(); // identifier
+    if (!StateInfo) {
+      Diag(StateLoc, diag::err_pragma_prefetch_invalid_arg);
+      while (Tok.isNot(tok::eof))
+        ConsumeAnyToken();
+      ConsumeToken();
+      return false;
+    }
+    Hint.StateLoc = IdentifierLoc::create(Actions.Context, StateLoc, StateInfo);
+
+    Name.setIdentifier(StateInfo, StateLoc);
+    ExprResult Res =
+        Actions.ActOnIdExpression(getCurScope(), ScopeSpec, TemplateKWLoc, Name,
+                                  Tok.is(tok::l_paren), false);
+
+    auto ConsumeRemainToks = [&]() {
+      while (Tok.isNot(tok::eof))
+        ConsumeAnyToken();
+      ConsumeToken();
+    };
+
+    if (Res.get()->containsErrors()) {
+      // Do not do error diagnosis here, it will be dealt with later.
+      ConsumeRemainToks();
+      return false;
+    }
+
+    if (OptionNoPrefetch) {
+      if (Toks.size() > 2) {
+        Diag(Tok.getLocation(), diag::warn_pragma_extra_tokens_at_eol)
+            << PragmaLoopHintString(Info->PragmaName, Info->Option);
+        while (Tok.isNot(tok::eof))
+          ConsumeAnyToken();
+      }
+      Hint.DeclRef = Res.get();
+    }
+
+    if (OptionPrefetch) {
+      if (Tok.isNot(tok::eof)) {
+        if (Tok.isNot(tok::comma)) {
+          Diag(Tok.getLocation(), diag::err_pragma_prefetch_punc)
+              << /*IsIdentifier=*/1 << StateInfo->getName();
+          ConsumeRemainToks();
+          return false;
+        }
+
+        PP.Lex(Tok);                               // ','
+        ExprResult R1 = ParseConstantExpression(); // 'numeric_constant'
+        if (R1.isInvalid() ||
+            Actions.CheckLoopHintExpr(R1.get(), R1.get()->getExprLoc(),
+                                      /*IsPrefetch=*/true)) {
+          ConsumeRemainToks();
+          return false;
+        }
+
+        if (Tok.isNot(tok::eof)) {
+          if (Tok.isNot(tok::comma)) {
+            Diag(Tok.getLocation(), diag::err_pragma_prefetch_punc)
+                << /*IsIdentifier=*/0;
+            ConsumeRemainToks();
+            return false;
+          }
+
+          PP.Lex(Tok);                               // ','
+          ExprResult R2 = ParseConstantExpression(); // 'numeric_constant'
+          if (R2.isInvalid() ||
+              Actions.CheckLoopHintExpr(R2.get(), R2.get()->getExprLoc())) {
+            ConsumeRemainToks();
+            return false;
+          }
+          Hint.ValueExpr1 = R2.get();
+
+          if (Tok.isNot(tok::eof)) {
+            Diag(Tok.getLocation(), diag::warn_pragma_extra_tokens_at_eol)
+                << PragmaLoopHintString(Info->PragmaName, Info->Option);
+            while (Tok.isNot(tok::eof))
+              ConsumeAnyToken();
+          }
+        }
+        Hint.ValueExpr = R1.get();
+      }
+      Hint.DeclRef = Res.get();
+    }
+
+    ConsumeToken(); // Consume the constant expression eof terminator.
   } else if (OptionInfo && OptionInfo->getName() == "vectorize_width") {
     PP.EnterTokenStream(Toks, /*DisableMacroExpansion=*/false,
                         /*IsReinject=*/false);
     ConsumeAnnotationToken();
 
     SourceLocation StateLoc = Toks[0].getLocation();
     IdentifierInfo *StateInfo = Toks[0].getIdentifierInfo();
     StringRef IsScalableStr = StateInfo ? StateInfo->getName() : "";

 ///    'interleave' '(' loop-hint-keyword ')'
 ///    'unroll' '(' unroll-hint-keyword ')'
 ///    'vectorize_predicate' '(' loop-hint-keyword ')'
 ///    'vectorize_width' '(' loop-hint-value ')'
 ///    'interleave_count' '(' loop-hint-value ')'
 ///    'unroll_count' '(' loop-hint-value ')'
 ///    'pipeline' '(' disable ')'
 ///    'pipeline_initiation_interval' '(' loop-hint-value ')'
+///    'noprefetch' '(' loop-access-decl ')'
+///    'prefetch' '(' loop-access-decl ',' loop-hint-value ',' loop-hint-value ')'
 ///
 ///  loop-hint-keyword:
 ///    'enable'
 ///    'disable'
 ///    'assume_safety'
 ///
 ///  unroll-hint-keyword:
 ///    'enable'
 ///    'disable'
 ///    'full'
 ///
 ///  loop-hint-value:
 ///    constant-expression
 ///
+///  loop-access-decl:
+///    declaration-reference-expression
+///
 /// Specifying vectorize(enable) or vectorize_width(_value_) instructs llvm to
 /// try vectorizing the instructions of the loop it precedes. Specifying
 /// interleave(enable) or interleave_count(_value_) instructs llvm to try
 /// interleaving multiple iterations of the loop it precedes. The width of the
 /// vector instructions is specified by vectorize_width() and the number of
 /// interleaved loop iterations is specified by interleave_count(). Specifying a
 /// value of 1 effectively disables vectorization/interleaving, even if it is
 /// possible and profitable, and 0 is invalid. The loop vectorizer currently
bool OptionValid = llvm::StringSwitch<bool>(OptionInfo->getName())
                            .Case("unroll", true)
                            .Case("distribute", true)
                            .Case("vectorize_predicate", true)
                            .Case("vectorize_width", true)
                            .Case("interleave_count", true)
                            .Case("unroll_count", true)
                            .Case("pipeline", true)
                            .Case("pipeline_initiation_interval", true)
+                           .Case("noprefetch", true)
+                           .Case("prefetch", true)
                            .Default(false);
     if (!OptionValid) {
       PP.Diag(Tok.getLocation(), diag::err_pragma_loop_invalid_option)
           << /*MissingOption=*/false << OptionInfo;
       return;
     }
     PP.Lex(Tok);
 

clang/lib/Parse/ParseStmt.cpp

StmtResult Parser::ParsePragmaLoopHint(StmtVector &Stmts,
   SourceLocation StartLoc = Tok.getLocation();
 
   // Get loop hints and consume annotated token.
   while (Tok.is(tok::annot_pragma_loop_hint)) {
     LoopHint Hint;
     if (!HandlePragmaLoopHint(Hint))
       continue;
 
-    ArgsUnion ArgHints[] = {Hint.PragmaNameLoc, Hint.OptionLoc, Hint.StateLoc,
-                            ArgsUnion(Hint.ValueExpr)};
+    ArgsUnion ArgHints[] = {Hint.PragmaNameLoc,
+                            Hint.OptionLoc,
+                            Hint.StateLoc,
+                            ArgsUnion(Hint.DeclRef),
+                            ArgsUnion(Hint.ValueExpr),
+                            ArgsUnion(Hint.ValueExpr1)};
     TempAttrs.addNew(Hint.PragmaNameLoc->Ident, Hint.Range, nullptr,
-                     Hint.PragmaNameLoc->Loc, ArgHints, 4,
+                     Hint.PragmaNameLoc->Loc, ArgHints, 6,
                      ParsedAttr::AS_Pragma);
   }
 
   // Get the next statement.
   MaybeParseCXX11Attributes(Attrs);
 
   ParsedAttributes EmptyDeclSpecAttrs(AttrFactory);
   StmtResult S = ParseStatementOrDeclarationAfterAttributes(

clang/lib/Sema/SemaExpr.cpp

S.Diag(Loc, diagnostic)
       << Ty
       << StringRef(buffer.data(), buffer.size());
   }
 
   bool isExact = (result == APFloat::opOK);
   return FloatingLiteral::Create(S.Context, Val, isExact, Ty, Loc);
 }
 
-bool Sema::CheckLoopHintExpr(Expr *E, SourceLocation Loc) {
+bool Sema::CheckLoopHintExpr(Expr *E, SourceLocation Loc, bool IsPrefetch) {
   assert(E && "Invalid expression");
 
   if (E->isValueDependent())
     return false;
 
   QualType QT = E->getType();
   if (!QT->isIntegerType() || QT->isBooleanType() || QT->isCharType()) {
     Diag(E->getExprLoc(), diag::err_pragma_loop_invalid_argument_type) << QT;
     return true;
   }
 
   llvm::APSInt ValueAPS;
   ExprResult R = VerifyIntegerConstantExpression(E, &ValueAPS);
 
   if (R.isInvalid())
     return true;
 
+  if (IsPrefetch) {
+    auto Val = ValueAPS.getSExtValue();
+    if (Val < 0 || Val > 3) {
+      Diag(E->getExprLoc(), diag::err_pragma_loop_invalid_numeric)
+          << toString(ValueAPS, 10);
+      return true;
+    }
+  } else {
     bool ValueIsPositive = ValueAPS.isStrictlyPositive();
     if (!ValueIsPositive || ValueAPS.getActiveBits() > 31) {
       Diag(E->getExprLoc(), diag::err_pragma_loop_invalid_argument_value)
           << toString(ValueAPS, 10) << ValueIsPositive;
       return true;
     }
+  }
 
   return false;
 }
 
 ExprResult Sema::ActOnNumericConstant(const Token &Tok, Scope *UDLScope) {
   // Fast path for a single digit (which is quite common).  A single digit
   // cannot have a trigraph, escaped newline, radix prefix, or suffix.
   if (Tok.getLength() == 1) {

clang/lib/Sema/SemaOpenMP.cpp

StmtResult Sema::ActOnOpenMPUnrollDirective(ArrayRef<OMPClause *> Clauses,
   // the maximum trip count, which will also generate a remainder loop. Just
   // `unroll(enable)` (which could have been useful if the user has not
   // specified a concrete factor; even though the outer loop cannot be
   // influenced anymore, would avoid more code bloat than necessary) will refuse
   // the loop because "Won't unroll; remainder loop could not be generated when
   // assuming runtime trip count". Even if it did work, it must not choose a
   // larger unroll factor than the maximum loop length, or it would always just
   // execute the remainder loop.
-  LoopHintAttr *UnrollHintAttr =
-      LoopHintAttr::CreateImplicit(Context, LoopHintAttr::UnrollCount,
-                                   LoopHintAttr::Numeric, MakeFactorExpr());
+  LoopHintAttr *UnrollHintAttr = LoopHintAttr::CreateImplicit(
+      Context, LoopHintAttr::UnrollCount, LoopHintAttr::Numeric, nullptr,
+      MakeFactorExpr(), nullptr);
   AttributedStmt *InnerUnrolled =
       AttributedStmt::Create(Context, StartLoc, {UnrollHintAttr}, InnerFor);
 
   // Outer For init-statement: auto .unrolled.iv = 0
   AddInitializerToDecl(
       OuterIVDecl, ActOnIntegerConstant(LoopHelper.Init->getExprLoc(), 0).get(),
       /*DirectInit=*/false);
   StmtResult OuterInit = new (Context)

clang/lib/Sema/SemaStmtAttr.cpp

return ::new (S.Context) SuppressAttr(
       S.Context, A, DiagnosticIdentifiers.data(), DiagnosticIdentifiers.size());
 }
 
 static Attr *handleLoopHintAttr(Sema &S, Stmt *St, const ParsedAttr &A,
                                 SourceRange) {
   IdentifierLoc *PragmaNameLoc = A.getArgAsIdent(0);
   IdentifierLoc *OptionLoc = A.getArgAsIdent(1);
   IdentifierLoc *StateLoc = A.getArgAsIdent(2);
-  Expr *ValueExpr = A.getArgAsExpr(3);
+  Expr *DeclRef = A.getArgAsExpr(3);
+  Expr *ValueExpr = A.getArgAsExpr(4);
+  Expr *ValueExpr1 = A.getArgAsExpr(5);
 
   StringRef PragmaName =
       llvm::StringSwitch<StringRef>(PragmaNameLoc->Ident->getName())
           .Cases("unroll", "nounroll", "unroll_and_jam", "nounroll_and_jam",
                  PragmaNameLoc->Ident->getName())
           .Default("clang loop");
 
   // This could be handled automatically by adding a Subjects definition in
Option = llvm::StringSwitch<LoopHintAttr::OptionType>(
                  .Case("vectorize_predicate", LoopHintAttr::VectorizePredicate)
                  .Case("interleave_count", LoopHintAttr::InterleaveCount)
                  .Case("unroll", LoopHintAttr::Unroll)
                  .Case("unroll_count", LoopHintAttr::UnrollCount)
                  .Case("pipeline", LoopHintAttr::PipelineDisabled)
                  .Case("pipeline_initiation_interval",
                        LoopHintAttr::PipelineInitiationInterval)
                  .Case("distribute", LoopHintAttr::Distribute)
+                 .Case("noprefetch", LoopHintAttr::NoPrefetch)
+                 .Case("prefetch", LoopHintAttr::Prefetch)
                  .Default(LoopHintAttr::Vectorize);
     if (Option == LoopHintAttr::VectorizeWidth) {
       assert((ValueExpr || (StateLoc && StateLoc->Ident)) &&
              "Attribute must have a valid value expression or argument.");
       if (ValueExpr && S.CheckLoopHintExpr(ValueExpr, St->getBeginLoc()))
         return nullptr;
       if (StateLoc && StateLoc->Ident && StateLoc->Ident->isStr("scalable"))
         State = LoopHintAttr::ScalableWidth;
       else
         State = LoopHintAttr::FixedWidth;
     } else if (Option == LoopHintAttr::InterleaveCount ||
                Option == LoopHintAttr::UnrollCount ||
                Option == LoopHintAttr::PipelineInitiationInterval) {
       assert(ValueExpr && "Attribute must have a valid value expression.");
       if (S.CheckLoopHintExpr(ValueExpr, St->getBeginLoc()))
         return nullptr;
       State = LoopHintAttr::Numeric;
+    } else if (Option == LoopHintAttr::NoPrefetch ||
+               Option == LoopHintAttr::Prefetch) {
+      assert(DeclRef && "Attribute must hava a valid declaration reference.");
+      State = LoopHintAttr::Reference;
     } else if (Option == LoopHintAttr::Vectorize ||
                Option == LoopHintAttr::Interleave ||
                Option == LoopHintAttr::VectorizePredicate ||
                Option == LoopHintAttr::Unroll ||
                Option == LoopHintAttr::Distribute ||
                Option == LoopHintAttr::PipelineDisabled) {
       assert(StateLoc && StateLoc->Ident && "Loop hint must have an argument");
       if (StateLoc->Ident->isStr("disable"))
         State = LoopHintAttr::Disable;
       else if (StateLoc->Ident->isStr("assume_safety"))
         State = LoopHintAttr::AssumeSafety;
       else if (StateLoc->Ident->isStr("full"))
         State = LoopHintAttr::Full;
       else if (StateLoc->Ident->isStr("enable"))
         State = LoopHintAttr::Enable;
       else
         llvm_unreachable("bad loop hint argument");
     } else
       llvm_unreachable("bad loop hint");
   }
 
-  return LoopHintAttr::CreateImplicit(S.Context, Option, State, ValueExpr, A);
+  return LoopHintAttr::CreateImplicit(S.Context, Option, State, DeclRef,
+                                      ValueExpr, ValueExpr1, A);
 }
 
 namespace {
 class CallExprFinder : public ConstEvaluatedExprVisitor<CallExprFinder> {
   bool FoundAsmStmt = false;
   std::vector<const CallExpr *> CallExprs;
 
 public:
enum CategoryType {
     // transformation.
     Unroll,
     // The loop distribution transformation only has a state form that is
     // exposed by #pragma clang loop distribute (enable | disable).
     Distribute,
     // The vector predication only has a state form that is exposed by
     // #pragma clang loop vectorize_predicate (enable | disable).
     VectorizePredicate,
+    // For NoPrefetch, it only has one variant: a declaration-reference form.
+    NoPrefetch,
+    // For Prefetch, it has three variants: a declaration-reference form and
+    // two numeric forms. The first numeric form range from 0 to 3, indicating
+    // the cache level. The second numeric form is a positive integer,
+    // indicating the iteration ahead.
+    Prefetch,
     // This serves as a indicator to how many category are listed in this enum.
     NumberOfCategories
   };
   // The following array accumulates the hints encountered while iterating
   // through the attributes to check for compatibility.
   struct {
     const LoopHintAttr *StateAttr;
     const LoopHintAttr *NumericAttr;
+    const LoopHintAttr *ReferenceAttr;
   } HintAttrs[CategoryType::NumberOfCategories] = {};
+  llvm::DenseMap<ValueDecl *, const LoopHintAttr *> PrefetchPrevAttr;
 
   for (const auto *I : Attrs) {
     const LoopHintAttr *LH = dyn_cast<LoopHintAttr>(I);
 
     // Skip non loop hint attributes
     if (!LH)
       continue;
 
case LoopHintAttr::Distribute:
       break;
     case LoopHintAttr::PipelineDisabled:
     case LoopHintAttr::PipelineInitiationInterval:
       Category = Pipeline;
       break;
     case LoopHintAttr::VectorizePredicate:
       Category = VectorizePredicate;
       break;
+    case LoopHintAttr::Prefetch:
+      Category = Prefetch;
+      break;
+    case LoopHintAttr::NoPrefetch:
+      Category = NoPrefetch;
+      break;
     };
 
     assert(Category != NumberOfCategories && "Unhandled loop hint option");
     auto &CategoryState = HintAttrs[Category];
     const LoopHintAttr *PrevAttr;
     if (Option == LoopHintAttr::Vectorize ||
         Option == LoopHintAttr::Interleave || Option == LoopHintAttr::Unroll ||
         Option == LoopHintAttr::UnrollAndJam ||
         Option == LoopHintAttr::VectorizePredicate ||
         Option == LoopHintAttr::PipelineDisabled ||
         Option == LoopHintAttr::Distribute) {
       // Enable|Disable|AssumeSafety hint.  For example, vectorize(enable).
       PrevAttr = CategoryState.StateAttr;
       CategoryState.StateAttr = LH;
+    } else if (Option == LoopHintAttr::Prefetch ||
+               Option == LoopHintAttr::NoPrefetch) {
+      // Reference hint. For example, prefetch(a).
+      PrevAttr = nullptr;
+      CategoryState.ReferenceAttr = LH;
     } else {
       // Numeric hint.  For example, vectorize_width(8).
       PrevAttr = CategoryState.NumericAttr;
       CategoryState.NumericAttr = LH;
     }
 
     PrintingPolicy Policy(S.Context.getLangOpts());
     SourceLocation OptionLoc = LH->getRange().getBegin();
     if (PrevAttr)
       // Cannot specify same type of attribute twice.
       S.Diag(OptionLoc, diag::err_pragma_loop_compatibility)
           << /*Duplicate=*/true << PrevAttr->getDiagnosticName(Policy)
           << LH->getDiagnosticName(Policy);
 
+    // Cannot specify prefetch and/or noprefetch for a same declaration
+    // reference twice.
+    if (CategoryState.ReferenceAttr) {
+      if (auto DeclRef = dyn_cast<DeclRefExpr>(LH->getDeclRef())) {
+        auto VD = DeclRef->getDecl();
+        auto Iter = PrefetchPrevAttr.find(VD);
+        if (Iter != PrefetchPrevAttr.end()) {
+          S.Diag(OptionLoc, diag::err_pragma_loop_compatibility)
+              << /*Duplicate=*/false << Iter->second->getDiagnosticName(Policy)
+              << CategoryState.ReferenceAttr->getDiagnosticName(Policy);
+        }
+        PrefetchPrevAttr.insert({VD, LH});
+      }
+    }
+
     if (CategoryState.StateAttr && CategoryState.NumericAttr &&
         (Category == Unroll || Category == UnrollAndJam ||
          CategoryState.StateAttr->getState() == LoopHintAttr::Disable)) {
       // Disable hints are not compatible with numeric hints of the same
       // category.  As a special case, numeric unroll hints are also not
       // compatible with enable or full form of the unroll pragma because these
       // directives indicate full unrolling.
       S.Diag(OptionLoc, diag::err_pragma_loop_compatibility)

clang/lib/Sema/SemaTemplateInstantiate.cpp

TemplateInstantiator::TransformLoopHintAttr(const LoopHintAttr *LH) {
 
   // Generate error if there is a problem with the value.
   if (getSema().CheckLoopHintExpr(TransformedExpr, LH->getLocation()))
     return LH;
 
   // Create new LoopHintValueAttr with integral expression in place of the
   // non-type template parameter.
   return LoopHintAttr::CreateImplicit(getSema().Context, LH->getOption(),
-                                      LH->getState(), TransformedExpr, *LH);
+                                      LH->getState(), LH->getDeclRef(),
+                                      TransformedExpr, LH->getValue1(), *LH);
 }
 
 ExprResult TemplateInstantiator::transformNonTypeTemplateParmRef(
     Decl *AssociatedDecl, const NonTypeTemplateParmDecl *parm,
     SourceLocation loc, TemplateArgument arg,
     std::optional<unsigned> PackIndex) {
   ExprResult result;
 

clang/test/CodeGenCXX/pragma-prefetch-noprefetch.cpp

+// RUN: %clang_cc1 -std=c++11 -emit-llvm -o - %s | FileCheck -check-prefix=CHECK %s
+// RUN: %clang_cc1 -std=c++11 -fopenmp -emit-llvm -o - %s | FileCheck -check-prefix=CHECK-OMP %s
+ 
+// Verify noprefetch metadata is generated correctly.
+int noprefetch_test(int *a, int n, int &sum) {
+  #pragma clang loop noprefetch(a)
+  for (int i = 0; i < n; ++i) {
+// CHECK: [[TMP1:%.*]] = load i32, ptr [[ARRAYIDX1:%.*]], align [[ALIGN:.*]], !llvm.loop.prefetch ![[NOPREFETCH1:.*]]
+    sum += a[i];
+  }
+
+  int *b = a;
+  #pragma clang loop noprefetch(a)
+  #pragma clang loop noprefetch(b)
+  for (int i = 0; i < n; ++i) {
+// CHECK: [[TMP2:%.*]] = load i32, ptr [[ARRAYIDX2:%.*]], align [[ALIGN]], !llvm.loop.prefetch ![[NOPREFETCH2:.*]]
+// CHECK: [[TMP3:%.*]] = load i32, ptr [[ARRAYIDX3:%.*]], align [[ALIGN]], !llvm.loop.prefetch ![[NOPREFETCH3:.*]]
+    sum += a[i] + b[i];
+  }
+
+  #pragma clang loop noprefetch(a,)
+  for (int i = 0; i < n; ++i) {
+// CHECK: [[TMP4:%.*]] = load i32, ptr [[ARRAYIDX4:%.*]], align [[ALIGN]], !llvm.loop.prefetch ![[NOPREFETCH4:.*]]
+    sum += a[i];
+  }
+
+  #pragma clang loop vectorize(enable) noprefetch(a) unroll(full) distribute(enable)
+  for (int i = 0; i < n; ++i) {
+// CHECK: [[TMP5:%.*]] = load i32, ptr [[ARRAYIDX5:%.*]], align [[ALIGN]], !llvm.loop.prefetch ![[NOPREFETCH5:.*]]
+    sum += a[i];
+  }
+
+  #pragma clang loop noprefetch(a)
+  for (int i = 0; i < n; ++i) {
+// CHECK: [[TMP6:%.*]] = load i32, ptr [[ARRAYIDX6:%.*]], align [[ALIGN]], !llvm.loop.prefetch ![[NOPREFETCH6:.*]]
+    sum += a[i];
+    for (int j = 0; j < n; ++j) {
+// CHECK: [[TMP7:%.*]] = load i32, ptr [[ARRAYIDX7:%.*]], align [[ALIGN]], !llvm.loop.prefetch ![[NOPREFETCH6]]
+      sum += a[j];
+    }
+  }
+
+  int c[n];
+  #pragma clang loop noprefetch(c)
+  for (int i = 0; i < n; ++i) {
+// CHECK: [[TMP8:%.*]] = load i32, ptr [[ARRAYIDX8:%.*]], align [[ALIGN]], !llvm.loop.prefetch ![[NOPREFETCH7:.*]]
+    sum += c[i];
+  }
+  return sum;
+}
+
+// Verify while loop is recoginized after sequnence of pragma clang loop directives.
+int while_noprefetch_test(int *a, int n, int &sum) {
+  int i = 0;
+  #pragma clang loop vectorize(enable)
+  #pragma clang loop interleave_count(4)
+  #pragma clang loop noprefetch(a)
+  #pragma clang loop unroll(full)
+  #pragma clang loop distribute(enable)
+  while (i++ < n) {
+// CHECK: [[TMP:%.*]] = load i32, ptr [[ARRAYIDX:%.*]], align [[ALIGN]], !llvm.loop.prefetch ![[NOPREFETCH8:.*]]
+    sum += a[i];
+  }
+  return sum;
+}
+
+// Verify prefetch metadata is generated correctly.
+int prefetch_test(int *a, int n, int &sum) {
+  #pragma clang loop prefetch(a)
+  for (int i = 0; i < n; ++i) {
+// CHECK: [[TMP1:%.*]] = load i32, ptr [[ARRAYIDX1:%.*]], align [[ALIGN]], !llvm.loop.prefetch ![[PREFETCH1:.*]]
+    sum += a[i];
+  }
+
+  #pragma clang loop prefetch(a, 1)
+  for (int i = 0; i < n; ++i) {
+// CHECK: [[TMP2:%.*]] = load i32, ptr [[ARRAYIDX2:%.*]], align [[ALIGN]], !llvm.loop.prefetch ![[PREFETCH2:.*]]
+    sum += a[i];
+  }
+
+  #pragma clang loop prefetch(a, 1, 8)
+  for (int i = 0; i < n; ++i) {
+// CHECK: [[TMP3:%.*]] = load i32, ptr [[ARRAYIDX3:%.*]], align [[ALIGN]], !llvm.loop.prefetch ![[PREFETCH3:.*]]
+    sum += a[i];
+  }
+
+  #pragma clang loop prefetch(a, 1, 8,)
+  for (int i = 0; i < n; ++i) {
+// CHECK: [[TMP4:%.*]] = load i32, ptr [[ARRAYIDX4:%.*]], align [[ALIGN]], !llvm.loop.prefetch ![[PREFETCH4:.*]]
+    sum += a[i];
+  }
+
+  int *b = a;
+  #pragma clang loop prefetch(a, 1, 8)
+  #pragma clang loop prefetch(b, 1)
+  for (int i = 0; i < n; ++i) {
+// CHECK: [[TMP5:%.*]] = load i32, ptr [[ARRAYIDX5:%.*]], align [[ALIGN]], !llvm.loop.prefetch ![[PREFETCH5:.*]]
+// CHECK: [[TMP6:%.*]] = load i32, ptr [[ARRAYIDX6:%.*]], align [[ALIGN]], !llvm.loop.prefetch ![[PREFETCH6:.*]]
+    sum += a[i] + b[i];
+  }
+
+  #pragma clang loop prefetch(a)
+  for (int i = 0; i < n; ++i) {
+// CHECK: [[TMP7:%.*]] = load i32, ptr [[ARRAYIDX7:%.*]], align [[ALIGN]], !llvm.loop.prefetch ![[PREFETCH7:.*]]
+    sum += a[i];
+    for (int j = 0; j < n; ++j) {
+// CHECK: [[TMP8:%.*]] = load i32, ptr [[ARRAYIDX8:%.*]], align [[ALIGN]], !llvm.loop.prefetch ![[PREFETCH7]]
+      sum += a[j];
+    }
+  }
+
+  #pragma clang loop prefetch(a, 1)
+  for (int i = 0; i < n; ++i) {
+// CHECK: [[TMP9:%.*]] = load i32, ptr [[ARRAYIDX9:%.*]], align [[ALIGN]], !llvm.loop.prefetch ![[PREFETCH8:.*]]
+    sum += a[i];
+    #pragma clang loop prefetch(a, 2)
+    for (int j = 0; j < n; ++j) {
+// CHECK: [[TMP10:%.*]] = load i32, ptr [[ARRAYIDX10:%.*]], align [[ALIGN]], !llvm.loop.prefetch ![[PREFETCH9:.*]]
+      sum += a[j];
+    }
+  }
+
+  return sum;
+}
+
+// Verify while loop is recoginized after sequnence of pragma clang loop directives.
+int while_prefetch_test(int *a, int n, int &sum) {
+  int i = 0;
+  #pragma clang loop vectorize(enable)
+  #pragma clang loop interleave_count(4)
+  #pragma clang loop prefetch(a)
+  #pragma clang loop unroll(full)
+  #pragma clang loop distribute(enable)
+  while (i++ < n) {
+// CHECK: [[TMP1:%.*]] = load i32, ptr [[ARRAYIDX1:%.*]], align [[ALIGN]], !llvm.loop.prefetch ![[PREFETCH10:.*]]
+    sum += a[i];
+  }
+  return sum;
+}
+
+int prefetch_noprefetch_test(int *a, int n, int &sum) {
+  #pragma clang loop noprefetch(a)
+  for (int i = 0; i < n; ++i) {
+// CHECK: [[TMP1:%.*]] = load i32, ptr [[ARRAYIDX1:%.*]], align [[ALIGN]], !llvm.loop.prefetch ![[NOPREFETCH9:.*]]
+    sum += a[i];
+    #pragma clang loop prefetch(a)
+    for (int j = 0; j < n; ++j) {
+// CHECK: [[TMP2:%.*]] = load i32, ptr [[ARRAYIDX2:%.*]], align [[ALIGN]], !llvm.loop.prefetch ![[PREFETCH11:.*]]
+      sum += a[j];
+    }
+  }
+
+  #pragma clang loop prefetch(a)
+  for (int i = 0; i < n; ++i) {
+// CHECK: [[TMP3:%.*]] = load i32, ptr [[ARRAYIDX3:%.*]], align [[ALIGN]], !llvm.loop.prefetch ![[PREFETCH12:.*]]
+    sum += a[i];
+    #pragma clang loop noprefetch(a)
+    for (int j = 0; j < n; ++j) {
+// CHECK: [[TMP4:%.*]] = load i32, ptr [[ARRAYIDX4:%.*]], align [[ALIGN]], !llvm.loop.prefetch ![[NOPREFETCH10:.*]]
+      sum += a[j];
+    }
+  }
+  return sum;
+}
+
+int openmp_prefetch_test(int *a, int n, int &sum) {
+  #pragma omp parallel for simd
+  #pragma clang loop prefetch(a, 0, 4)
+  for (int i = 0; i < n; ++i) {
+// CHECK-OMP: [[TMP1:%.*]] = load i32, ptr [[ARRARYIDX1:%.*]], align [[ALIGN:.*]], !llvm.access.group [[GROUP:.*]], !llvm.loop.prefetch ![[PREFETCH13:.*]]
+    sum += a[i];
+  }
+
+  #pragma omp parallel for private(a)
+  #pragma clang loop prefetch(a, 1, 4)
+  for (int i = 0; i < n; ++i) {
+// CHECK-OMP: [[TMP2:%.*]] = load i32, ptr [[ARRARYIDX2:%.*]], align [[ALIGN]], !llvm.loop.prefetch ![[PREFETCH14:.*]]
+    sum += a[i];
+  }
+
+  #pragma omp parallel for firstprivate(a)
+  #pragma clang loop prefetch(a, 2, 4)
+  for (int i = 0; i < n; ++i) {
+// CHECK-OMP: [[TMP3:%.*]] = load i32, ptr [[ARRARYIDX3:%.*]], align [[ALIGN]], !llvm.loop.prefetch ![[PREFETCH15:.*]]
+    sum += a[i];
+  }
+  return sum;
+}
+
+// CHECK: ![[NOPREFETCH1]] = distinct !{i1 false, i32 -1, i32 -1}
+// CHECK: ![[NOPREFETCH2]] = distinct !{i1 false, i32 -1, i32 -1}
+// CHECK: ![[NOPREFETCH3]] = distinct !{i1 false, i32 -1, i32 -1}
+// CHECK: ![[NOPREFETCH4]] = distinct !{i1 false, i32 -1, i32 -1}
+// CHECK: ![[NOPREFETCH5]] = distinct !{i1 false, i32 -1, i32 -1}
+// CHECK: ![[NOPREFETCH6]] = distinct !{i1 false, i32 -1, i32 -1}
+// CHECK: ![[NOPREFETCH7]] = distinct !{i1 false, i32 -1, i32 -1}
+// CHECK: ![[NOPREFETCH8]] = distinct !{i1 false, i32 -1, i32 -1}
+// CHECK: ![[PREFETCH1]] = distinct !{i1 true, i32 -1, i32 -1}
+// CHECK: ![[PREFETCH2]] = distinct !{i1 true, i32 1, i32 -1}
+// CHECK: ![[PREFETCH3]] = distinct !{i1 true, i32 1, i32 8}
+// CHECK: ![[PREFETCH4]] = distinct !{i1 true, i32 1, i32 8}
+// CHECK: ![[PREFETCH5]] = distinct !{i1 true, i32 1, i32 8}
+// CHECK: ![[PREFETCH6]] = distinct !{i1 true, i32 1, i32 -1}
+// CHECK: ![[PREFETCH7]] = distinct !{i1 true, i32 -1, i32 -1}
+// CHECK: ![[PREFETCH8]] = distinct !{i1 true, i32 1, i32 -1}
+// CHECK: ![[PREFETCH9]] = distinct !{i1 true, i32 2, i32 -1}
+// CHECK: ![[PREFETCH10]] = distinct !{i1 true, i32 -1, i32 -1}
+// CHECK: ![[NOPREFETCH9]] = distinct !{i1 false, i32 -1, i32 -1}
+// CHECK: ![[PREFETCH11]] = distinct !{i1 true, i32 -1, i32 -1}
+// CHECK: ![[PREFETCH12]] = distinct !{i1 true, i32 -1, i32 -1}
+// CHECK: ![[NOPREFETCH10]] = distinct !{i1 false, i32 -1, i32 -1}
+// CHECK-OMP: ![[PREFETCH13]] = distinct !{i1 true, i32 0, i32 4}
+// CHECK-OMP: ![[PREFETCH14]] = distinct !{i1 true, i32 1, i32 4}
+// CHECK-OMP: ![[PREFETCH15]] = distinct !{i1 true, i32 2, i32 4}

clang/test/Parser/pragma-loop.cpp

 /* expected-error {{expected ')'}} */ #pragma clang loop interleave_count(4
 /* expected-error {{expected ')'}} */ #pragma clang loop unroll_count(4
 
 /* expected-error {{missing argument; expected 'enable', 'assume_safety' or 'disable'}} */ #pragma clang loop vectorize()
 /* expected-error {{missing argument; expected an integer value}} */ #pragma clang loop interleave_count()
 /* expected-error {{missing argument; expected 'enable', 'full' or 'disable'}} */ #pragma clang loop unroll()
 /* expected-error {{missing argument; expected 'enable' or 'disable'}} */ #pragma clang loop distribute()
 
-/* expected-error {{missing option; expected vectorize, vectorize_width, interleave, interleave_count, unroll, unroll_count, pipeline, pipeline_initiation_interval, vectorize_predicate, or distribute}} */ #pragma clang loop
+/* expected-error {{missing option; expected vectorize, vectorize_width, interleave, interleave_count, unroll, unroll_count, pipeline, pipeline_initiation_interval, vectorize_predicate, distribute, prefetch, or noprefetch}} */ #pragma clang loop
 /* expected-error {{invalid option 'badkeyword'}} */ #pragma clang loop badkeyword
 /* expected-error {{invalid option 'badkeyword'}} */ #pragma clang loop badkeyword(enable)
 /* expected-error {{invalid option 'badkeyword'}} */ #pragma clang loop vectorize(enable) badkeyword(4)
 /* expected-warning {{extra tokens at end of '#pragma clang loop'}} */ #pragma clang loop vectorize(enable) ,
   while (i-4 < Length) {
     List[i] = i;
   }
 

clang/test/Parser/pragma-prefetch-noprefetch.cpp

+// RUN: %clang_cc1 -std=c++11 -verify %s
+
+// Note that this puts the expected lines before the directives to work around
+// limitations in the -verify mode.
+
+int test_noprefetch(int *a, int n, int &sum) {
+  #pragma clang loop noprefetch(a)
+  for (int i = 0; i < n; ++i) {
+    sum += a[i];
+  }
+ 
+  int i = 0;
+  #pragma clang loop noprefetch(a)
+  while (i++ < n) {
+    sum += a[i];
+  }
+
+  #pragma clang loop noprefetch(a)
+/* expected-error {{incompatible directives 'noprefetch(a)' and 'prefetch(a)'}} */ #pragma clang loop prefetch(a)
+  for (int i = 0; i < n; ++i) {
+    sum += a[i];
+  }
+
+  #pragma clang loop noprefetch(a)
+/* expected-error {{incompatible directives 'noprefetch(a)' and 'noprefetch(a)'}} */ #pragma clang loop noprefetch(a)
+  for (int i = 0; i < n; ++i) {
+    sum += a[i];
+  }
+
+/* expected-error {{use of undeclared identifier 'b'}} */ #pragma clang loop noprefetch(b)
+  for (int i = 0; i < n; ++i) {
+    sum += a[i];
+  }
+
+/* expected-error {{use of undeclared identifier 'b'}} */ #pragma clang loop noprefetch(b)
+  for (int i = 0; i < n; ++i) {
+  int *b = a;
+    sum += a[i];
+  }
+
+/* expected-error {{invalid argument; expected an declared variable}} */ #pragma clang loop noprefetch(1)
+  for (int i = 0; i < n; ++i) {
+    sum += a[i];
+  }
+
+/* expected-error {{missing argument; expected an declared variable}} */ #pragma clang loop noprefetch()
+  for (int i = 0; i < n; ++i) {
+    sum += a[i];
+  }
+
+/* expected-warning {{extra tokens at end of '#pragma clang loop noprefetch' - ignored}} */ #pragma clang loop noprefetch(a,2)
+  for (int i = 0; i < n; ++i) {
+    sum += a[i];
+  }
+  return sum;
+}
+
+int test_prefetch(int *a, int n, int &sum) {
+  #pragma clang loop prefetch(a)
+  for (int i = 0; i < n; ++i) {
+    sum += a[i];
+  }
+
+  #pragma clang loop prefetch(a, 0)
+  for (int i = 0; i < n; ++i) {
+    sum += a[i];
+  }
+
+  #pragma clang loop prefetch(a, 1+2)
+  for (int i = 0; i < n; ++i) {
+    sum += a[i];
+  }
+
+  #pragma clang loop prefetch(a, 3, 2147483647)
+  for (int i = 0; i < n; ++i) {
+    sum += a[i];
+  }
+
+  int i = 0;
+  #pragma clang loop noprefetch(a)
+  while (i++ < n) {
+    sum += a[i];
+  }
+
+  int *b = a;
+  #pragma clang loop prefetch(a)
+  #pragma clang loop prefetch(b)
+  for (int i = 0; i < n; ++i) {
+    sum += a[i] + b[i];
+  }
+
+  #pragma clang loop prefetch(a)
+/* expected-error {{incompatible directives 'prefetch(a)' and 'prefetch(a, 1)'}} */ #pragma clang loop prefetch(a, 1)
+  for (int i = 0; i < n; ++i) {
+    sum += a[i];
+  }
+
+/* expected-error {{use of undeclared identifier 'c'}} */ #pragma clang loop prefetch(c)
+  for (int i = 0; i < n; ++i) {
+    sum += a[i];
+  }
+
+/* expected-error {{use of undeclared identifier 'c'}} */ #pragma clang loop prefetch(c)
+  for (int i = 0; i < n; ++i) {
+  int *c = a;
+    sum += a[i];
+  }
+
+/* expected-error {{invalid argument; expected an declared variable}} */ #pragma clang loop prefetch(1)
+  for (int i = 0; i < n; ++i) {
+    sum += a[i];
+  }
+
+/* expected-error {{missing argument; expected an declared variable}} */ #pragma clang loop prefetch()
+  for (int i = 0; i < n; ++i) {
+    sum += a[i];
+  }
+
+/* expected-error {{expected ')' or ',' after 'a'}} */ #pragma clang loop prefetch(a;)
+  for (int i = 0; i < n; ++i) {
+    sum += a[i];
+  }
+
+/* expected-error {{expected ')' or ','}} */ #pragma clang loop prefetch(a, 1;)
+  for (int i = 0; i < n; ++i) {
+    sum += a[i];
+  }
+
+/* expected-error {{expected expression}} */ #pragma clang loop prefetch(a,)
+  for (int i = 0; i < n; ++i) {
+    sum += a[i];
+  }
+
+/* expected-error {{invalid argument of type 'double'; expected an integer type}} */ #pragma clang loop prefetch(a, 2.5)
+  for (int i = 0; i < n; ++i) {
+    sum += a[i];
+  }
+
+/* expected-error {{expected expression}} */ #pragma clang loop prefetch(a,,)
+  for (int i = 0; i < n; ++i) {
+    sum += a[i];
+  }
+
+/* expected-error {{invalid argument of type 'char'; expected an integer type}} */ #pragma clang loop prefetch(a, 'c')
+  for (int i = 0; i < n; ++i) {
+    sum += a[i];
+  }
+
+/* expected-error {{use of undeclared identifier 'c'}} */ #pragma clang loop prefetch(a, c)
+  for (int i = 0; i < n; ++i) {
+    sum += a[i];
+  }
+
+/* expected-error {{invalid argument of type 'int *'; expected an integer type}} */ #pragma clang loop prefetch(a, a)
+  for (int i = 0; i < n; ++i) {
+    sum += a[i];
+  }
+
+/* expected-error {{invalid value '-1'; expected integer constant between 0 and 3}} */ #pragma clang loop prefetch(a, -1)
+  for (int i = 0; i < n; ++i) {
+    sum += a[i];
+  }
+
+/* expected-error {{invalid value '4'; expected integer constant between 0 and 3}} */ #pragma clang loop prefetch(a, 4)
+  for (int i = 0; i < n; ++i) {
+    sum += a[i];
+  }
+
+/* expected-error {{invalid value '5'; expected integer constant between 0 and 3}} */ #pragma clang loop prefetch(a, 1+4)
+  for (int i = 0; i < n; ++i) {
+    sum += a[i];
+  }
+
+/* expected-error {{expected expression}} */ #pragma clang loop prefetch(a, 2,)
+  for (int i = 0; i < n; ++i) {
+    sum += a[i];
+  }
+
+/* expected-error {{invalid argument of type 'char'; expected an integer type}} */ #pragma clang loop prefetch(a, 2, 'c')
+  for (int i = 0; i < n; ++i) {
+    sum += a[i];
+  }
+
+/* expected-error {{invalid value '0'; must be positive}} */ #pragma clang loop prefetch(a, 2, 0)
+  for (int i = 0; i < n; ++i) {
+    sum += a[i];
+  }
+
+/* expected-error {{value '2147483648' is too large}} */ #pragma clang loop prefetch(a, 2, 2147483648)
+  for (int i = 0; i < n; ++i) {
+    sum += a[i];
+  }
+
+/* expected-warning {{extra tokens at end of '#pragma clang loop noprefetch' - ignored}} */ #pragma clang loop noprefetch(a,2,3,4)
+  for (int i = 0; i < n; ++i) {
+    sum += a[i];
+  }
+
+  return sum;
+}

clang/test/Parser/pragma-unroll-and-jam.cpp

 #pragma unroll(4)
   for (int i = 0; i < Length; i++) {
     for (int j = 0; j < Length; j++) {
       List[i * Length + j] = Value;
     }
   }
 
 // pragma clang unroll_and_jam is disabled for the moment
-/* expected-error {{invalid option 'unroll_and_jam'; expected vectorize, vectorize_width, interleave, interleave_count, unroll, unroll_count, pipeline, pipeline_initiation_interval, vectorize_predicate, or distribute}} */ #pragma clang loop unroll_and_jam(4)
+/* expected-error {{invalid option 'unroll_and_jam'; expected vectorize, vectorize_width, interleave, interleave_count, unroll, unroll_count, pipeline, pipeline_initiation_interval, vectorize_predicate, distribute, prefetch, or noprefetch}} */ #pragma clang loop unroll_and_jam(4)
   for (int i = 0; i < Length; i++) {
     for (int j = 0; j < Length; j++) {
       List[i * Length + j] = Value;
     }
   }
 
 #pragma unroll_and_jam
 /* expected-error {{expected statement}} */ }