Merging r322003:

------------------------------------------------------------------------
r322003 | niravd | 2018-01-08 08:21:35 -0800 (Mon, 08 Jan 2018) | 11 lines

[DAG] Teach BaseIndexOffset to correctly handle with indexed operations

BaseIndexOffset address analysis incorrectly ignores offsets folded
into indexed memory operations causing potential errors in alias
analysis of pre-indexed operations.

Reviewers: efriedma, RKSimon, hfinkel, jyknight

Subscribers: hiraditya, javed.absar, llvm-commits

Differential Revision: https://reviews.llvm.org/D41701
------------------------------------------------------------------------

llvm-svn: 322693

README.md

@@ -1,5 +0,0 @@
-# Low Level Virtual Machine (LLVM)
-
-This directory and its subdirectories contain source code for LLVM,
-a toolkit for the construction of highly optimized compilers,
-optimizers, and runtime environments.

clang-tools-extra/clang-tidy/fuchsia/OverloadedOperatorCheck.cpp

@@ -30,8 +30,12 @@ void OverloadedOperatorCheck::registerMatchers(MatchFinder *Finder) {
 }
 
 void OverloadedOperatorCheck::check(const MatchFinder::MatchResult &Result) {
-  if (const auto *D = Result.Nodes.getNodeAs<FunctionDecl>("decl"))
-    diag(D->getLocStart(), "cannot overload %0") << D;
+  const auto *D = Result.Nodes.getNodeAs<FunctionDecl>("decl");
+  assert(D && "No FunctionDecl captured!");
+  
+  SourceLocation Loc = D->getLocStart();
+  if (Loc.isValid())
+    diag(Loc, "cannot overload %0") << D;
 }
 
 } // namespace fuchsia

clang-tools-extra/test/clang-tidy/fuchsia-overloaded-operator.cpp

@@ -16,3 +16,6 @@ public:
 
 A operator-(const A &A1, const A &A2);
 // CHECK-MESSAGES: [[@LINE-1]]:1: warning: cannot overload 'operator-' [fuchsia-overloaded-operator]
+
+void operator delete(void*, void*) throw();
+// CHECK-MESSAGES: [[@LINE-1]]:1: warning: cannot overload 'operator delete' [fuchsia-overloaded-operator]

clang/docs/ClangCommandLineReference.rst

@@ -120,6 +120,10 @@ Output path for the plist report
 
 .. option:: -compatibility\_version<arg>
 
+.. option:: --config <arg>
+
+Specifies configuration file
+
 .. option:: --constant-cfstrings
 
 .. option:: -coverage, --coverage
@@ -1545,6 +1549,10 @@ Enable ARC-style weak references in Objective-C
 
 OpenMP target code is compiled as relocatable using the -c flag. For OpenMP targets the code is relocatable by default.
 
+.. option:: -fopenmp-simd, -fno-openmp-simd
+
+Emit OpenMP code only for SIMD-based constructs.
+
 .. option:: -fopenmp-use-tls
 
 .. option:: -fopenmp-version=<arg>
@@ -1998,7 +2006,7 @@ Link stack frames through backchain on System Z
 
 .. option:: -mconsole<arg>
 
-.. option:: -mcpu=<arg>, -mv4 (equivalent to -mcpu=hexagonv4), -mv5 (equivalent to -mcpu=hexagonv5), -mv55 (equivalent to -mcpu=hexagonv55), -mv60 (equivalent to -mcpu=hexagonv60), -mv62 (equivalent to -mcpu=hexagonv62)
+.. option:: -mcpu=<arg>, -mv4 (equivalent to -mcpu=hexagonv4), -mv5 (equivalent to -mcpu=hexagonv5), -mv55 (equivalent to -mcpu=hexagonv55), -mv60 (equivalent to -mcpu=hexagonv60), -mv62 (equivalent to -mcpu=hexagonv62), -mv65 (equivalent to -mcpu=hexagonv65)
 
 .. option:: -mdefault-build-attributes<arg>, -mno-default-build-attributes<arg>
 
@@ -2328,6 +2336,8 @@ X86
 
 .. option:: -mavx2, -mno-avx2
 
+.. option:: -mavx512bitalg, -mno-avx512bitalg
+
 .. option:: -mavx512bw, -mno-avx512bw
 
 .. option:: -mavx512cd, -mno-avx512cd
@@ -2344,8 +2354,12 @@ X86
 
 .. option:: -mavx512vbmi, -mno-avx512vbmi
 
+.. option:: -mavx512vbmi2, -mno-avx512vbmi2
+
 .. option:: -mavx512vl, -mno-avx512vl
 
+.. option:: -mavx512vnni, -mno-avx512vnni
+
 .. option:: -mavx512vpopcntdq, -mno-avx512vpopcntdq
 
 .. option:: -mbmi, -mno-bmi
@@ -2370,6 +2384,8 @@ X86
 
 .. option:: -mfxsr, -mno-fxsr
 
+.. option:: -mgfni, -mno-gfni
+
 .. option:: -mibt, -mno-ibt
 
 .. option:: -mlwp, -mno-lwp
@@ -2424,6 +2440,10 @@ X86
 
 .. option:: -mtbm, -mno-tbm
 
+.. option:: -mvaes, -mno-vaes
+
+.. option:: -mvpclmulqdq, -mno-vpclmulqdq
+
 .. option:: -mx87, -m80387, -mno-x87
 
 .. option:: -mxop, -mno-xop

clang/docs/OpenMPSupport.rst

@@ -0,0 +1,68 @@
+.. raw:: html
+
+  <style type="text/css">
+    .none { background-color: #FFCCCC }
+    .partial { background-color: #FFFF99 }
+    .good { background-color: #CCFF99 }
+  </style>
+
+.. role:: none
+.. role:: partial
+.. role:: good
+
+==================
+OpenMP Support
+==================
+
+Clang fully supports OpenMP 3.1 + some elements of OpenMP 4.5. Clang supports offloading to X86_64, AArch64 and PPC64[LE] devices.
+Support for Cuda devices is not ready yet.
+The status of major OpenMP 4.5 features support in Clang.
+
+Standalone directives
+=====================
+
+* #pragma omp [for] simd: :good:`Complete`.
+
+* #pragma omp declare simd: :partial:`Partial`.  We support parsing/semantic
+  analysis + generation of special attributes for X86 target, but still
+  missing the LLVM pass for vectorization.
+
+* #pragma omp taskloop [simd]: :good:`Complete`.
+
+* #pragma omp target [enter|exit] data: :good:`Complete`.
+
+* #pragma omp target update: :good:`Complete`.
+
+* #pragma omp target: :partial:`Partial`.  No support for the `depend` clauses.
+
+* #pragma omp declare target: :partial:`Partial`.  No full codegen support.
+
+* #pragma omp teams: :good:`Complete`.
+
+* #pragma omp distribute [simd]: :good:`Complete`.
+
+* #pragma omp distribute parallel for [simd]: :good:`Complete`.
+
+Combined directives
+===================
+
+* #pragma omp parallel for simd: :good:`Complete`.
+
+* #pragma omp target parallel: :partial:`Partial`.  No support for the `depend` clauses.
+
+* #pragma omp target parallel for [simd]: :partial:`Partial`.  No support for the `depend` clauses.
+
+* #pragma omp target simd: :partial:`Partial`.  No support for the `depend` clauses.
+
+* #pragma omp target teams: :partial:`Partial`.  No support for the `depend` clauses.
+
+* #pragma omp teams distribute [simd]: :good:`Complete`.
+
+* #pragma omp target teams distribute [simd]: :partial:`Partial`.  No support for the and `depend` clauses.
+
+* #pragma omp teams distribute parallel for [simd]: :good:`Complete`.
+
+* #pragma omp target teams distribute parallel for [simd]: :partial:`Partial`.  No full codegen support.
+
+Clang does not support any constructs/updates from upcoming OpenMP 5.0 except for `reduction`-based clauses in the `task` and `target`-based directives.
+

clang/docs/ReleaseNotes.rst

@@ -168,8 +168,8 @@ Attribute Changes in Clang
 Windows Support
 ---------------
 
-Clang's support for building native Windows programs ...
-
+- Clang now has initial, preliminary support for targeting Windows on
+  ARM64.
 
 C Language Changes in Clang
 ---------------------------

clang/docs/index.rst

@@ -39,6 +39,7 @@ Using Clang as a Compiler
    SourceBasedCodeCoverage
    Modules
    MSVCCompatibility
+   OpenMPSupport
    ThinLTO
    CommandGuide/index
    FAQ

clang/include/clang/Basic/Attr.td

@@ -549,6 +549,7 @@ def Aligned : InheritableAttr {
                                           Keyword<"_Alignas">]>,
                    Accessor<"isDeclspec",[Declspec<"align">]>];
   let Documentation = [Undocumented];
+  let DuplicatesAllowedWhileMerging = 1;
 }
 
 def AlignValue : Attr {

clang/include/clang/Basic/BuiltinsX86.def

@@ -1357,15 +1357,15 @@ TARGET_BUILTIN(__builtin_ia32_vpshrdvw128_maskz, "V8sV8sV8sV8sUc", "", "avx512vl
 TARGET_BUILTIN(__builtin_ia32_vpshrdvw256_maskz, "V16sV16sV16sV16sUs", "", "avx512vl,avx512vbmi2")
 TARGET_BUILTIN(__builtin_ia32_vpshrdvw512_maskz, "V32sV32sV32sV32sUi", "", "avx512vbmi2")
 
-TARGET_BUILTIN(__builtin_ia32_vpshrdd128_mask, "V4iV4iV4iiV4iUc", "", "avx512vl,avx512vbmi2")
-TARGET_BUILTIN(__builtin_ia32_vpshrdd256_mask, "V8iV8iV8iiV8iUc", "", "avx512vl,avx512vbmi2")
-TARGET_BUILTIN(__builtin_ia32_vpshrdd512_mask, "V16iV16iV16iiV16iUs", "", "avx512vbmi2")
-TARGET_BUILTIN(__builtin_ia32_vpshrdq128_mask, "V2LLiV2LLiV2LLiiV2LLiUc", "", "avx512vl,avx512vbmi2")
-TARGET_BUILTIN(__builtin_ia32_vpshrdq256_mask, "V4LLiV4LLiV4LLiiV4LLiUc", "", "avx512vl,avx512vbmi2")
-TARGET_BUILTIN(__builtin_ia32_vpshrdq512_mask, "V8LLiV8LLiV8LLiiV8LLiUc", "", "avx512vbmi2")
-TARGET_BUILTIN(__builtin_ia32_vpshrdw128_mask, "V8sV8sV8siV8sUc", "", "avx512vl,avx512vbmi2")
-TARGET_BUILTIN(__builtin_ia32_vpshrdw256_mask, "V16sV16sV16siV16sUs", "", "avx512vl,avx512vbmi2")
-TARGET_BUILTIN(__builtin_ia32_vpshrdw512_mask, "V32sV32sV32siV32sUi", "", "avx512vbmi2")
+TARGET_BUILTIN(__builtin_ia32_vpshrdd128_mask, "V4iV4iV4iIiV4iUc", "", "avx512vl,avx512vbmi2")
+TARGET_BUILTIN(__builtin_ia32_vpshrdd256_mask, "V8iV8iV8iIiV8iUc", "", "avx512vl,avx512vbmi2")
+TARGET_BUILTIN(__builtin_ia32_vpshrdd512_mask, "V16iV16iV16iIiV16iUs", "", "avx512vbmi2")
+TARGET_BUILTIN(__builtin_ia32_vpshrdq128_mask, "V2LLiV2LLiV2LLiIiV2LLiUc", "", "avx512vl,avx512vbmi2")
+TARGET_BUILTIN(__builtin_ia32_vpshrdq256_mask, "V4LLiV4LLiV4LLiIiV4LLiUc", "", "avx512vl,avx512vbmi2")
+TARGET_BUILTIN(__builtin_ia32_vpshrdq512_mask, "V8LLiV8LLiV8LLiIiV8LLiUc", "", "avx512vbmi2")
+TARGET_BUILTIN(__builtin_ia32_vpshrdw128_mask, "V8sV8sV8sIiV8sUc", "", "avx512vl,avx512vbmi2")
+TARGET_BUILTIN(__builtin_ia32_vpshrdw256_mask, "V16sV16sV16sIiV16sUs", "", "avx512vl,avx512vbmi2")
+TARGET_BUILTIN(__builtin_ia32_vpshrdw512_mask, "V32sV32sV32sIiV32sUi", "", "avx512vbmi2")
 
 TARGET_BUILTIN(__builtin_ia32_pmovswb512_mask, "V32cV32sV32cUi", "", "avx512bw")
 TARGET_BUILTIN(__builtin_ia32_pmovuswb512_mask, "V32cV32sV32cUi", "", "avx512bw")

clang/include/clang/Basic/TokenKinds.def

@@ -398,7 +398,6 @@ TYPE_TRAIT_2(__builtin_types_compatible_p, TypeCompatible, KEYNOCXX)
 KEYWORD(__builtin_va_arg            , KEYALL)
 KEYWORD(__extension__               , KEYALL)
 KEYWORD(__float128                  , KEYALL)
-ALIAS("_Float128", __float128       , KEYNOCXX)
 KEYWORD(__imag                      , KEYALL)
 KEYWORD(__int128                    , KEYALL)
 KEYWORD(__label__                   , KEYALL)

clang/include/clang/StaticAnalyzer/Core/BugReporter/BugType.h

@@ -32,27 +32,39 @@ private:
   const CheckName Check;
   const std::string Name;
   const std::string Category;
-  bool SuppressonSink;
+  const CheckerBase *Checker;
+  bool SuppressOnSink;
 
   virtual void anchor();
-public:
-  BugType(class CheckName check, StringRef name, StringRef cat)
-      : Check(check), Name(name), Category(cat), SuppressonSink(false) {}
-  BugType(const CheckerBase *checker, StringRef name, StringRef cat)
-      : Check(checker->getCheckName()), Name(name), Category(cat),
-        SuppressonSink(false) {}
-  virtual ~BugType() {}
 
-  // FIXME: Should these be made strings as well?
+public:
+  BugType(CheckName Check, StringRef Name, StringRef Cat)
+      : Check(Check), Name(Name), Category(Cat), Checker(nullptr),
+        SuppressOnSink(false) {}
+  BugType(const CheckerBase *Checker, StringRef Name, StringRef Cat)
+      : Check(Checker->getCheckName()), Name(Name), Category(Cat),
+        Checker(Checker), SuppressOnSink(false) {}
+  virtual ~BugType() = default;
+
   StringRef getName() const { return Name; }
   StringRef getCategory() const { return Category; }
-  StringRef getCheckName() const { return Check.getName(); }
+  StringRef getCheckName() const {
+    // FIXME: This is a workaround to ensure that the correct check name is used
+    // The check names are set after the constructors are run.
+    // In case the BugType object is initialized in the checker's ctor
+    // the Check field will be empty. To circumvent this problem we use 
+    // CheckerBase whenever it is possible.
+    StringRef CheckName =
+        Checker ? Checker->getCheckName().getName() : Check.getName();
+    assert(!CheckName.empty() && "Check name is not set properly.");
+    return CheckName;
+  }
 
   /// isSuppressOnSink - Returns true if bug reports associated with this bug
   ///  type should be suppressed if the end node of the report is post-dominated
   ///  by a sink node.
-  bool isSuppressOnSink() const { return SuppressonSink; }
-  void setSuppressOnSink(bool x) { SuppressonSink = x; }
+  bool isSuppressOnSink() const { return SuppressOnSink; }
+  void setSuppressOnSink(bool x) { SuppressOnSink = x; }
 
   virtual void FlushReports(BugReporter& BR);
 };
@@ -74,7 +86,7 @@ public:
   StringRef getDescription() const { return desc; }
 };
 
-} // end GR namespace
+} // end ento namespace
 
 } // end clang namespace
 #endif

clang/lib/AST/ODRHash.cpp

@@ -478,6 +478,8 @@ void ODRHash::AddFunctionDecl(const FunctionDecl *Function) {
 
   // TODO: Fix hashing for class methods.
   if (isa<CXXMethodDecl>(Function)) return;
+  // And friend functions.
+  if (Function->getFriendObjectKind()) return;
 
   // Skip functions that are specializations or in specialization context.
   const DeclContext *DC = Function;

clang/lib/CodeGen/CGBuiltin.cpp

@@ -915,7 +915,11 @@ EmitCheckedMixedSignMultiply(CodeGenFunction &CGF, const clang::Expr *Op1,
       Overflow = CGF.Builder.CreateOr(Overflow, TruncOverflow);
     }
 
-    Result = CGF.Builder.CreateTrunc(UnsignedResult, ResTy);
+    // Negate the product if it would be negative in infinite precision.
+    Result = CGF.Builder.CreateSelect(
+        IsNegative, CGF.Builder.CreateNeg(UnsignedResult), UnsignedResult);
+
+    Result = CGF.Builder.CreateTrunc(Result, ResTy);
   }
   assert(Overflow && Result && "Missing overflow or result");
 

clang/lib/CodeGen/ObjectFilePCHContainerOperations.cpp

@@ -229,6 +229,11 @@ public:
       Builder->getModuleDebugInfo()->completeRequiredType(RD);
   }
 
+  void HandleImplicitImportDecl(ImportDecl *D) override {
+    if (!D->getImportedOwningModule())
+      Builder->getModuleDebugInfo()->EmitImportDecl(*D);
+  }
+
   /// Emit a container holding the serialized AST.
   void HandleTranslationUnit(ASTContext &Ctx) override {
     assert(M && VMContext && Builder);

clang/lib/Frontend/InitPreprocessor.cpp

@@ -817,10 +817,6 @@ static void InitializePredefinedMacros(const TargetInfo &TI,
   DefineFloatMacros(Builder, "FLT", &TI.getFloatFormat(), "F");
   DefineFloatMacros(Builder, "DBL", &TI.getDoubleFormat(), "");
   DefineFloatMacros(Builder, "LDBL", &TI.getLongDoubleFormat(), "L");
-  if (TI.hasFloat128Type())
-    // FIXME: Switch away from the non-standard "Q" when we can
-    DefineFloatMacros(Builder, "FLT128", &TI.getFloat128Format(), "Q");
-
 
   // Define a __POINTER_WIDTH__ macro for stdint.h.
   Builder.defineMacro("__POINTER_WIDTH__",

clang/lib/Lex/Lexer.cpp

@@ -2009,18 +2009,21 @@ bool Lexer::LexAngledStringLiteral(Token &Result, const char *CurPtr) {
   const char *AfterLessPos = CurPtr;
   char C = getAndAdvanceChar(CurPtr, Result);
   while (C != '>') {
-    // Skip escaped characters.
-    if (C == '\\' && CurPtr < BufferEnd) {
-      // Skip the escaped character.
-      getAndAdvanceChar(CurPtr, Result);
-    } else if (C == '\n' || C == '\r' ||             // Newline.
-               (C == 0 && (CurPtr-1 == BufferEnd ||  // End of file.
-                           isCodeCompletionPoint(CurPtr-1)))) {
+    // Skip escaped characters.  Escaped newlines will already be processed by
+    // getAndAdvanceChar.
+    if (C == '\\')
+      C = getAndAdvanceChar(CurPtr, Result);
+
+    if (C == '\n' || C == '\r' ||             // Newline.
+        (C == 0 && (CurPtr-1 == BufferEnd ||  // End of file.
+                    isCodeCompletionPoint(CurPtr-1)))) {
       // If the filename is unterminated, then it must just be a lone <
       // character.  Return this as such.
       FormTokenWithChars(Result, AfterLessPos, tok::less);
       return true;
-    } else if (C == 0) {
+    }
+
+    if (C == 0) {
       NulCharacter = CurPtr-1;
     }
     C = getAndAdvanceChar(CurPtr, Result);

clang/lib/Sema/SemaTemplateDeduction.cpp

@@ -502,6 +502,10 @@ DeduceTemplateArguments(Sema &S,
                         SmallVectorImpl<DeducedTemplateArgument> &Deduced) {
   assert(Arg.isCanonical() && "Argument type must be canonical");
 
+  // Treat an injected-class-name as its underlying template-id.
+  if (auto *Injected = dyn_cast<InjectedClassNameType>(Arg))
+    Arg = Injected->getInjectedSpecializationType();
+
   // Check whether the template argument is a dependent template-id.
   if (const TemplateSpecializationType *SpecArg
         = dyn_cast<TemplateSpecializationType>(Arg)) {

clang/lib/Sema/SemaTemplateInstantiateDecl.cpp

@@ -4160,7 +4160,8 @@ void Sema::BuildVariableInstantiation(
   // it right away if the type contains 'auto'.
   if ((!isa<VarTemplateSpecializationDecl>(NewVar) &&
        !InstantiatingVarTemplate &&
-       !(OldVar->isInline() && OldVar->isThisDeclarationADefinition())) ||
+       !(OldVar->isInline() && OldVar->isThisDeclarationADefinition() &&
+         !NewVar->isThisDeclarationADefinition())) ||
       NewVar->getType()->isUndeducedType())
     InstantiateVariableInitializer(NewVar, OldVar, TemplateArgs);
 

clang/lib/StaticAnalyzer/Checkers/MallocChecker.cpp

@@ -2900,8 +2900,13 @@ void ento::registerNewDeleteLeaksChecker(CheckerManager &mgr) {
       mgr.getCurrentCheckName();
   // We currently treat NewDeleteLeaks checker as a subchecker of NewDelete
   // checker.
-  if (!checker->ChecksEnabled[MallocChecker::CK_NewDeleteChecker])
+  if (!checker->ChecksEnabled[MallocChecker::CK_NewDeleteChecker]) {
     checker->ChecksEnabled[MallocChecker::CK_NewDeleteChecker] = true;
+    // FIXME: This does not set the correct name, but without this workaround
+    //        no name will be set at all.
+    checker->CheckNames[MallocChecker::CK_NewDeleteChecker] =
+        mgr.getCurrentCheckName();
+  }
 }
 
 #define REGISTER_CHECKER(name)                                                 \

clang/lib/StaticAnalyzer/Checkers/ValistChecker.cpp

@@ -64,7 +64,7 @@ private:
                                  CheckerContext &C) const;
   void reportLeakedVALists(const RegionVector &LeakedVALists, StringRef Msg1,
                            StringRef Msg2, CheckerContext &C, ExplodedNode *N,
-                           bool ForceReport = false) const;
+                           bool ReportUninit = false) const;
 
   void checkVAListStartCall(const CallEvent &Call, CheckerContext &C,
                             bool IsCopy) const;
@@ -267,15 +267,19 @@ void ValistChecker::reportUninitializedAccess(const MemRegion *VAList,
 void ValistChecker::reportLeakedVALists(const RegionVector &LeakedVALists,
                                         StringRef Msg1, StringRef Msg2,
                                         CheckerContext &C, ExplodedNode *N,
-                                        bool ForceReport) const {
+                                        bool ReportUninit) const {
   if (!(ChecksEnabled[CK_Unterminated] ||
-        (ChecksEnabled[CK_Uninitialized] && ForceReport)))
+        (ChecksEnabled[CK_Uninitialized] && ReportUninit)))
     return;
   for (auto Reg : LeakedVALists) {
     if (!BT_leakedvalist) {
-      BT_leakedvalist.reset(new BugType(CheckNames[CK_Unterminated],
-                                        "Leaked va_list",
-                                        categories::MemoryError));
+      // FIXME: maybe creating a new check name for this type of bug is a better
+      // solution.
+      BT_leakedvalist.reset(
+          new BugType(CheckNames[CK_Unterminated].getName().empty()
+                          ? CheckNames[CK_Uninitialized]
+                          : CheckNames[CK_Unterminated],
+                      "Leaked va_list", categories::MemoryError));
       BT_leakedvalist->setSuppressOnSink(true);
     }
 
@@ -375,7 +379,7 @@ void ValistChecker::checkVAListEndCall(const CallEvent &Call,
 
 std::shared_ptr<PathDiagnosticPiece> ValistChecker::ValistBugVisitor::VisitNode(
     const ExplodedNode *N, const ExplodedNode *PrevN, BugReporterContext &BRC,
-    BugReport &BR) {
+    BugReport &) {
   ProgramStateRef State = N->getState();
   ProgramStateRef StatePrev = PrevN->getState();
 

clang/test/Analysis/malloc.c

@@ -1720,13 +1720,6 @@ void *smallocWarn(size_t size) {
   }
 }
 
-char *dupstrWarn(const char *s) {
-  const int len = strlen(s);
-  char *p = (char*) smallocWarn(len + 1);
-  strcpy(p, s); // expected-warning{{String copy function overflows destination buffer}}
-  return p;
-}
-
 int *radar15580979() {
   int *data = (int *)malloc(32);
   int *p = data ?: (int*)malloc(32); // no warning

clang/test/CodeGen/builtins-overflow.c

@@ -373,7 +373,9 @@ int test_mixed_sign_mull_overflow_unsigned(int x, unsigned y) {
 // CHECK-NEXT:  [[NotNull:%.*]] = icmp ne i32 [[UnsignedResult]], 0
 // CHECK-NEXT:  [[Underflow:%.*]] = and i1 [[IsNeg]], [[NotNull]]
 // CHECK-NEXT:  [[OFlow:%.*]] = or i1 [[UnsignedOFlow]], [[Underflow]]
-// CHECK-NEXT:  store i32 [[UnsignedResult]], i32* %{{.*}}, align 4
+// CHECK-NEXT:  [[NegatedResult:%.*]] = sub i32 0, [[UnsignedResult]]
+// CHECK-NEXT:  [[Result:%.*]] = select i1 [[IsNeg]], i32 [[NegatedResult]], i32 [[UnsignedResult]]
+// CHECK-NEXT:  store i32 [[Result]], i32* %{{.*}}, align 4
 // CHECK:       br i1 [[OFlow]]
 
   unsigned result;
@@ -432,7 +434,9 @@ long long test_mixed_sign_mulll_overflow_trunc_unsigned(long long x, unsigned lo
 // CHECK-NEXT:  [[OVERFLOW_PRE_TRUNC:%.*]] = or i1 {{.*}}, [[UNDERFLOW]]
 // CHECK-NEXT:  [[TRUNC_OVERFLOW:%.*]] = icmp ugt i64 [[UNSIGNED_RESULT]], 4294967295
 // CHECK-NEXT:  [[OVERFLOW:%.*]] = or i1 [[OVERFLOW_PRE_TRUNC]], [[TRUNC_OVERFLOW]]
-// CHECK-NEXT:  trunc i64 [[UNSIGNED_RESULT]] to i32
+// CHECK-NEXT:  [[NEGATED:%.*]] = sub i64 0, [[UNSIGNED_RESULT]]
+// CHECK-NEXT:  [[RESULT:%.*]] = select i1 {{.*}}, i64 [[NEGATED]], i64 [[UNSIGNED_RESULT]]
+// CHECK-NEXT:  trunc i64 [[RESULT]] to i32
 // CHECK-NEXT:  store
   unsigned result;
   if (__builtin_mul_overflow(y, x, &result))

clang/test/CodeGenCXX/cxx1z-inline-variables.cpp

@@ -58,14 +58,22 @@ template<typename T> struct X {
   static int a;
   static inline int b;
   static int c;
+  static const int d;
+  static int e;
 };
 // CHECK: @_ZN1XIiE1aE = linkonce_odr global i32 10
 // CHECK: @_ZN1XIiE1bE = global i32 20
 // CHECK-NOT: @_ZN1XIiE1cE
+// CHECK: @_ZN1XIiE1dE = linkonce_odr constant i32 40
+// CHECK: @_ZN1XIiE1eE = linkonce_odr global i32 50
 template<> inline int X<int>::a = 10;
 int &use3 = X<int>::a;
 template<> int X<int>::b = 20;
 template<> inline int X<int>::c = 30;
+template<typename T> constexpr int X<T>::d = 40;
+template<typename T> inline int X<T>::e = 50;
+const int *use_x_int_d = &X<int>::d;
+const int *use_x_int_e = &X<int>::e;
 
 template<typename T> struct Y;
 template<> struct Y<int> {

clang/test/Modules/ExtDebugInfo.cpp

@@ -187,7 +187,7 @@ void foo() {
 
 // CHECK: !DIGlobalVariable(name: "anon_enum", {{.*}}, type: ![[ANON_ENUM:[0-9]+]]
 // CHECK: !DICompositeType(tag: DW_TAG_enumeration_type, scope: ![[NS]],
-// CHECK-SAME:             line: 16
+// CHECK-SAME:             line: 19
 
 // CHECK: !DIGlobalVariable(name: "GlobalUnion",
 // CHECK-SAME:              type: ![[GLOBAL_UNION:[0-9]+]]

clang/test/Modules/Inputs/DebugCXX.h

@@ -1,4 +1,7 @@
 /* -*- C++ -*- */
+
+#include "dummy.h"
+
 namespace DebugCXX {
   // Records.
   struct Struct {

clang/test/Modules/Inputs/odr_hash-Friend/Box.h

@@ -0,0 +1,14 @@
+template <class T>
+struct iterator {
+  void Compare(const iterator &x) { }
+  friend void Check(iterator) {}
+};
+
+template <class T = int> struct Box {
+  iterator<T> I;
+
+  void test() {
+    Check(I);
+    I.Compare(I);
+  }
+};

clang/test/Modules/Inputs/odr_hash-Friend/M1.h

@@ -0,0 +1,6 @@
+#include "Box.h"
+
+void Peek() {
+  Box<> Gift;
+  Gift.test();
+}

clang/test/Modules/Inputs/odr_hash-Friend/M2.h

@@ -0,0 +1,5 @@
+#include "Box.h"
+void x() {
+  Box<> Unused;
+  //Unused.test();
+}

clang/test/Modules/Inputs/odr_hash-Friend/M3.h

@@ -0,0 +1,7 @@
+#include "Box.h"
+#include "M2.h"
+
+void Party() {
+  Box<> Present;
+  Present.test();
+}

clang/test/Modules/Inputs/odr_hash-Friend/module.modulemap

@@ -0,0 +1,15 @@
+module Box {
+  header "Box.h"
+}
+
+module Module1 {
+  header "M1.h"
+}
+
+module Module2 {
+  header "M2.h"
+}
+
+module Module3 {
+  header "M3.h"
+}

clang/test/Modules/ModuleDebugInfo.cpp

@@ -5,12 +5,13 @@
 
 // Modules:
 // RUN: rm -rf %t
-// RUN: %clang_cc1 -triple %itanium_abi_triple -x objective-c++ -std=c++11 -debug-info-kind=limited -fmodules -fmodule-format=obj -fimplicit-module-maps -DMODULES -fmodules-cache-path=%t %s -I %S/Inputs -I %t -emit-llvm -o %t.ll -mllvm -debug-only=pchcontainer &>%t-mod.ll
+// RUN: %clang_cc1 -triple %itanium_abi_triple -x objective-c++ -std=c++11 -debugger-tuning=lldb -debug-info-kind=limited -fmodules -fmodule-format=obj -fimplicit-module-maps -DMODULES -fmodules-cache-path=%t %s -I %S/Inputs -I %t -emit-llvm -o %t.ll -mllvm -debug-only=pchcontainer &>%t-mod.ll
 // RUN: cat %t-mod.ll | FileCheck %s
 // RUN: cat %t-mod.ll | FileCheck --check-prefix=CHECK-NEG %s
+// RUN: cat %t-mod.ll | FileCheck --check-prefix=CHECK-MOD %s
 
 // PCH:
-// RUN: %clang_cc1 -triple %itanium_abi_triple -x c++ -std=c++11 -emit-pch -fmodule-format=obj -I %S/Inputs -o %t.pch %S/Inputs/DebugCXX.h -mllvm -debug-only=pchcontainer &>%t-pch.ll
+// RUN: %clang_cc1 -triple %itanium_abi_triple -x c++ -std=c++11  -debugger-tuning=lldb -emit-pch -fmodule-format=obj -I %S/Inputs -o %t.pch %S/Inputs/DebugCXX.h -mllvm -debug-only=pchcontainer &>%t-pch.ll
 // RUN: cat %t-pch.ll | FileCheck %s
 // RUN: cat %t-pch.ll | FileCheck --check-prefix=CHECK-NEG %s
 
@@ -18,6 +19,9 @@
 @import DebugCXX;
 #endif
 
+// CHECK-MOD: distinct !DICompileUnit(language: DW_LANG_{{.*}}C_plus_plus,
+// CHECK-MOD: distinct !DICompileUnit(language: DW_LANG_{{.*}}C_plus_plus,
+
 // CHECK: distinct !DICompileUnit(language: DW_LANG_{{.*}}C_plus_plus,
 // CHECK-SAME:                    isOptimized: false,
 // CHECK-NOT:                     splitDebugFilename:
@@ -27,6 +31,8 @@
 // CHECK-SAME:             identifier: "_ZTSN8DebugCXX4EnumE")
 // CHECK: !DINamespace(name: "DebugCXX"
 
+// CHECK-MOD: ![[DEBUGCXX:.*]] = !DIModule(scope: null, name: "DebugCXX
+
 // CHECK: !DICompositeType(tag: DW_TAG_enumeration_type,
 // CHECK-NOT:              name:
 // CHECK-SAME:             )
@@ -150,4 +156,11 @@
 // CHECK-SAME:                             name: "WithSpecializedBase<float>",
 // CHECK-SAME:                             flags: DIFlagFwdDecl,
 
+// CHECK-MOD: !DIImportedEntity(tag: DW_TAG_imported_declaration, scope: ![[DEBUGCXX]],
+// CHECK-MOD-SAME:              entity: ![[DUMMY:[0-9]+]],
+// CHECK-MOD-SAME:              line: 3)
+// CHECK-MOD: ![[DUMMY]] = !DIModule(scope: null, name: "dummy",
+// CHECK-MOD: distinct !DICompileUnit(language: DW_LANG_ObjC_plus_plus,
+// CHECK-MOD-SAME:  splitDebugFilename: "{{.*}}dummy{{.*}}.pcm",
+
 // CHECK-NEG-NOT: !DICompositeType(tag: DW_TAG_structure_type, name: "PureForwardDecl"

clang/test/Modules/odr_hash-Friend.cpp

@@ -0,0 +1,22 @@
+// RUN: rm -rf %t
+
+// RUN: %clang_cc1 -fmodules -fmodules-cache-path=%t/modules.cache \
+// RUN:  -I %S/Inputs/odr_hash-Friend \
+// RUN:  -emit-obj -o /dev/null \
+// RUN:  -fmodules \
+// RUN:  -fimplicit-module-maps \
+// RUN:  -fmodules-cache-path=%t/modules.cache \
+// RUN:  -std=c++11 -x c++ %s -verify
+
+// PR35939: MicrosoftMangle.cpp triggers an assertion failure on this test.
+// UNSUPPORTED: system-windows
+
+// expected-no-diagnostics
+
+#include "Box.h"
+#include "M1.h"
+#include "M3.h"
+
+void Run() {
+  Box<> Present;
+}

clang/test/Preprocessor/cuda-types.cu

@@ -9,40 +9,40 @@
 
 // RUN: %clang --cuda-host-only -nocudainc -target i386-unknown-linux-gnu -x cuda -E -dM -o - /dev/null \
 // RUN:   | grep 'define __[^ ]*\(TYPE\|MAX\|SIZEOF|WIDTH\)\|define __GCC_ATOMIC' \
-// RUN:   | grep -v '__FLT128\|__LDBL\|_LONG_DOUBLE' > %t/i386-host-defines-filtered
+// RUN:   | grep -v '__LDBL\|_LONG_DOUBLE' > %t/i386-host-defines-filtered
 // RUN: %clang --cuda-device-only -nocudainc -nocudalib -target i386-unknown-linux-gnu -x cuda -E -dM -o - /dev/null \
 // RUN:   | grep 'define __[^ ]*\(TYPE\|MAX\|SIZEOF|WIDTH\)\|define __GCC_ATOMIC' \
-// RUN:   | grep -v '__FLT128\|__LDBL\|_LONG_DOUBLE' > %t/i386-device-defines-filtered
+// RUN:   | grep -v '__LDBL\|_LONG_DOUBLE' > %t/i386-device-defines-filtered
 // RUN: diff %t/i386-host-defines-filtered %t/i386-device-defines-filtered
 
 // RUN: %clang --cuda-host-only -nocudainc -target x86_64-unknown-linux-gnu -x cuda -E -dM -o - /dev/null \
 // RUN:   | grep 'define __[^ ]*\(TYPE\|MAX\|SIZEOF|WIDTH\)\|define __GCC_ATOMIC' \
-// RUN:   | grep -v '__FLT128\|__LDBL\|_LONG_DOUBLE' > %t/x86_64-host-defines-filtered
+// RUN:   | grep -v '__LDBL\|_LONG_DOUBLE' > %t/x86_64-host-defines-filtered
 // RUN: %clang --cuda-device-only -nocudainc -nocudalib -target x86_64-unknown-linux-gnu -x cuda -E -dM -o - /dev/null \
 // RUN:   | grep 'define __[^ ]*\(TYPE\|MAX\|SIZEOF|WIDTH\)\|define __GCC_ATOMIC' \
-// RUN:   | grep -v '__FLT128\|__LDBL\|_LONG_DOUBLE' > %t/x86_64-device-defines-filtered
+// RUN:   | grep -v '__LDBL\|_LONG_DOUBLE' > %t/x86_64-device-defines-filtered
 // RUN: diff %t/x86_64-host-defines-filtered %t/x86_64-device-defines-filtered
 
 // RUN: %clang --cuda-host-only -nocudainc -target powerpc64-unknown-linux-gnu -x cuda -E -dM -o - /dev/null \
 // RUN:   | grep 'define __[^ ]*\(TYPE\|MAX\|SIZEOF|WIDTH\)\|define __GCC_ATOMIC' \
-// RUN:   | grep -v '__FLT128\|__LDBL\|_LONG_DOUBLE' > %t/powerpc64-host-defines-filtered
+// RUN:   | grep -v '__LDBL\|_LONG_DOUBLE' > %t/powerpc64-host-defines-filtered
 // RUN: %clang --cuda-device-only -nocudainc -nocudalib -target powerpc64-unknown-linux-gnu -x cuda -E -dM -o - /dev/null \
 // RUN:   | grep 'define __[^ ]*\(TYPE\|MAX\|SIZEOF|WIDTH\)\|define __GCC_ATOMIC' \
-// RUN:   | grep -v '__FLT128\|__LDBL\|_LONG_DOUBLE' > %t/powerpc64-device-defines-filtered
+// RUN:   | grep -v '__LDBL\|_LONG_DOUBLE' > %t/powerpc64-device-defines-filtered
 // RUN: diff %t/powerpc64-host-defines-filtered %t/powerpc64-device-defines-filtered
 
 // RUN: %clang --cuda-host-only -nocudainc -target i386-windows-msvc -x cuda -E -dM -o - /dev/null \
 // RUN:   | grep 'define __[^ ]*\(TYPE\|MAX\|SIZEOF|WIDTH\)\|define __GCC_ATOMIC' \
-// RUN:   | grep -v '__FLT128\|__LDBL\|_LONG_DOUBLE' > %t/i386-msvc-host-defines-filtered
+// RUN:   | grep -v '__LDBL\|_LONG_DOUBLE' > %t/i386-msvc-host-defines-filtered
 // RUN: %clang --cuda-device-only -nocudainc -nocudalib -target i386-windows-msvc -x cuda -E -dM -o - /dev/null \
 // RUN:   | grep 'define __[^ ]*\(TYPE\|MAX\|SIZEOF|WIDTH\)\|define __GCC_ATOMIC' \
-// RUN:   | grep -v '__FLT128\|__LDBL\|_LONG_DOUBLE' > %t/i386-msvc-device-defines-filtered
+// RUN:   | grep -v '__LDBL\|_LONG_DOUBLE' > %t/i386-msvc-device-defines-filtered
 // RUN: diff %t/i386-msvc-host-defines-filtered %t/i386-msvc-device-defines-filtered
 
 // RUN: %clang --cuda-host-only -nocudainc -target x86_64-windows-msvc -x cuda -E -dM -o - /dev/null \
 // RUN:   | grep 'define __[^ ]*\(TYPE\|MAX\|SIZEOF|WIDTH\)\|define __GCC_ATOMIC' \
-// RUN:   | grep -v '__FLT128\|__LDBL\|_LONG_DOUBLE' > %t/x86_64-msvc-host-defines-filtered
+// RUN:   | grep -v '__LDBL\|_LONG_DOUBLE' > %t/x86_64-msvc-host-defines-filtered
 // RUN: %clang --cuda-device-only -nocudainc -nocudalib -target x86_64-windows-msvc -x cuda -E -dM -o - /dev/null \
 // RUN:   | grep 'define __[^ ]*\(TYPE\|MAX\|SIZEOF|WIDTH\)\|define __GCC_ATOMIC' \
-// RUN:   | grep -v '__FLT128\|__LDBL\|_LONG_DOUBLE' > %t/x86_64-msvc-device-defines-filtered
+// RUN:   | grep -v '__LDBL\|_LONG_DOUBLE' > %t/x86_64-msvc-device-defines-filtered
 // RUN: diff %t/x86_64-msvc-host-defines-filtered %t/x86_64-msvc-device-defines-filtered

clang/test/Sema/_Float128.c

@@ -1,22 +0,0 @@
-// RUN: %clang_cc1 -verify %s
-// RUN: %clang_cc1 -triple powerpc64-linux -verify %s
-// RUN: %clang_cc1 -triple i686-windows-gnu -verify %s
-// RUN: %clang_cc1 -triple x86_64-windows-gnu -verify %s
-// RUN: %clang_cc1 -triple x86_64-windows-msvc -verify %s
-
-#if defined(__FLOAT128__) || defined(__SIZEOF_FLOAT128__)
-_Float128 f;
-_Float128 tiny = __FLT128_EPSILON__;
-int g(int x, _Float128 *y) {
-  return x + *y;
-}
-
-// expected-no-diagnostics
-#else
-_Float128 f;  // expected-error {{__float128 is not supported on this target}}
-float tiny = __FLT128_EPSILON__; // expected-error{{use of undeclared identifier}}
-int g(int x, _Float128 *y) {  // expected-error {{__float128 is not supported on this target}}
-  return x + *y;
-}
-
-#endif  // defined(__FLOAT128__) || defined(__SIZEOF_FLOAT128__)

clang/test/SemaCXX/cxx1z-class-template-argument-deduction.cpp

@@ -309,6 +309,17 @@ namespace dependent {
   template int New(int);
 }
 
+namespace injected_class_name {
+  template<typename T = void> struct A {
+    A();
+    template<typename U> A(A<U>);
+  };
+  A<int> a;
+  A b = a;
+  using T = decltype(a);
+  using T = decltype(b);
+}
+
 #else
 
 // expected-no-diagnostics

clang/test/SemaTemplate/alignas.cpp

@@ -21,3 +21,14 @@ struct C { char a[16]; };
 
 static_assert(sizeof(my_union<A, B, C>) == 16, "");
 static_assert(alignof(my_union<A, B, C>) == 8, "");
+
+namespace PR35028 {
+  template<class X, int Alignment> struct alignas(X) alignas(long long) alignas(long double) alignas(Alignment) Aligned {
+    union {
+      long long align1;
+      long double align2;
+      char data[sizeof(X)];
+    };
+  };
+  Aligned<int, 1> a;
+}

clang/test/SemaTemplate/cxx17-inline-variables.cpp

@@ -16,3 +16,14 @@ namespace CompleteType {
 
   constexpr int n = X<true>::value;
 }
+
+template <typename T> struct A {
+  static const int n;
+  static const int m;
+  constexpr int f() { return n; }
+  constexpr int g() { return n; }
+};
+template <typename T> constexpr int A<T>::n = sizeof(A) + sizeof(T);
+template <typename T> inline constexpr int A<T>::m = sizeof(A) + sizeof(T);
+static_assert(A<int>().f() == 5);
+static_assert(A<int>().g() == 5);

clang/unittests/Lex/LexerTest.cpp

@@ -476,6 +476,8 @@ TEST_F(LexerTest, GetBeginningOfTokenWithEscapedNewLine) {
 TEST_F(LexerTest, AvoidPastEndOfStringDereference) {
   std::vector<Token> LexedTokens = Lex("  //  \\\n");
   EXPECT_TRUE(LexedTokens.empty());
+  EXPECT_TRUE(Lex("#include <\\\\").empty());
+  EXPECT_TRUE(Lex("#include <\\\\\n").empty());
 }
 
 TEST_F(LexerTest, StringizingRasString) {

debuginfo-tests/.arcconfig

@@ -1,3 +0,0 @@
-{
-  "conduit_uri" : "https://reviews.llvm.org/"
-}

debuginfo-tests/CMakeLists.txt

@@ -1,26 +0,0 @@
-# Debug Info tests.  These tests invoke clang to generate programs with
-# various types of debug info, and then run those programs under a debugger
-# such as GDB or LLDB to verify the results.
-
-set(DEBUGINFO_TESTS_SOURCE_DIR ${CMAKE_CURRENT_SOURCE_DIR})
-set(DEBUGINFO_TESTS_BINARY_DIR ${CMAKE_CURRENT_BINARY_DIR})
-
-set(DEBUGINFO_TEST_DEPS
-  clang
-  FileCheck
-  count
-  not
-  )
-
-configure_lit_site_cfg(
-  ${CMAKE_CURRENT_SOURCE_DIR}/lit.site.cfg.py.in
-  ${CMAKE_CURRENT_BINARY_DIR}/lit.site.cfg.py
-  MAIN_CONFIG
-  ${CMAKE_CURRENT_SOURCE_DIR}/lit.cfg.py
-  )
-
-add_lit_testsuite(check-debuginfo "Running debug info integration tests"
-  ${CMAKE_CURRENT_BINARY_DIR}
-  DEPENDS ${DEBUGINFO_TEST_DEPS}
-  )
-set_target_properties(check-debuginfo PROPERTIES FOLDER "Debug info tests")

debuginfo-tests/README.txt

@@ -1,19 +0,0 @@
-                                                                   -*- rst -*-
-This is a collection of tests to check debugging information generated by 
-compiler. This test suite can be checked out inside clang/test folder. This 
-will enable 'make test' for clang to pick up these tests. Typically, test 
-cases included here includes debugger commands and intended debugger output 
-as comments in source file using DEBUGGER: and CHECK: as prefixes respectively.
-
-For example::
-
-  define i32 @f1(i32 %i) nounwind ssp {
-  ; DEBUGGER: break f1
-  ; DEBUGGER: r
-  ; DEBUGGER: p i 
-  ; CHECK: $1 = 42 
-  entry:
-  }
-
-is a testcase where the debugger is asked to break at function 'f1' and 
-print value of argument 'i'. The expected value of 'i' is 42 in this case.

debuginfo-tests/aggregate-indirect-arg.cpp

@@ -1,32 +0,0 @@
-// RUN: %clangxx %target_itanium_abi_host_triple -O0 -g %s -c -o %t.o
-// RUN: %clangxx %target_itanium_abi_host_triple %t.o -o %t.out
-// RUN: %test_debuginfo %s %t.out 
-// Radar 8945514
-// DEBUGGER: break 22
-// DEBUGGER: r
-// DEBUGGER: p v
-// CHECK: ${{[0-9]+}} =
-// CHECK:  Data ={{.*}} 0x0{{(0*)}}
-// CHECK:  Kind = 2142
-
-class SVal {
-public:
-  ~SVal() {}
-  const void* Data;
-  unsigned Kind;
-};
-
-void bar(SVal &v) {}
-class A {
-public:
-  void foo(SVal v) { bar(v); }
-};
-
-int main() {
-  SVal v;
-  v.Data = 0;
-  v.Kind = 2142;
-  A a;
-  a.foo(v);
-  return 0;
-}

debuginfo-tests/asan-blocks.c

@@ -1,41 +0,0 @@
-// RUN: %clang -fblocks %target_itanium_abi_host_triple -arch x86_64 %s -o %t.out -g -fsanitize=address
-// RUN: %test_debuginfo %s %t.out
-// FIXME: Remove system-darwin when we build BlocksRuntime everywhere.
-// REQUIRES: not_asan, system-darwin
-//           Zorg configures the ASAN stage2 bots to not build the asan
-//           compiler-rt. Only run this test on non-asanified configurations.
-void b();
-struct S {
-  int a[8];
-};
-
-int f(struct S s, unsigned i) {
-  // DEBUGGER: break 17
-  // DEBUGGER: r
-  // DEBUGGER: p s
-  // CHECK: a = ([0] = 0, [1] = 1, [2] = 2, [3] = 3, [4] = 4, [5] = 5, [6] = 6, [7] = 7)
-  return s.a[i];
-}
-
-int main(int argc, const char **argv) {
-  struct S s = {{0, 1, 2, 3, 4, 5, 6, 7}};
-  if (f(s, 4) == 4) {
-    // DEBUGGER: break 27
-    // DEBUGGER: c
-    // DEBUGGER: p s
-    // CHECK: a = ([0] = 0, [1] = 1, [2] = 2, [3] = 3, [4] = 4, [5] = 5, [6] = 6, [7] = 7)
-    b();
-  }
-  return 0;
-}
-
-void c() {}
-
-void b() {
-  // DEBUGGER: break 40
-  // DEBUGGER: c
-  // DEBUGGER: p x
-  // CHECK: 42
-  __block int x = 42;
-  c();
-}

debuginfo-tests/asan.c

@@ -1,31 +0,0 @@
-// RUN: %clang -fblocks %target_itanium_abi_host_triple -arch x86_64 %s -o %t.out -g -fsanitize=address
-// RUN: %test_debuginfo %s %t.out
-// REQUIRES: not_asan
-//           Zorg configures the ASAN stage2 bots to not build the asan
-//           compiler-rt. Only run this test on non-asanified configurations.
-//
-
-struct S {
-  int a[8];
-};
-
-int f(struct S s, unsigned i) {
-  // DEBUGGER: break 14
-  return s.a[i];
-}
-
-int main(int argc, const char **argv) {
-  struct S s = {{0, 1, 2, 3, 4, 5, 6, 7}};
-  if (f(s, 4) == 4)
-    return f(s, 0);
-  return 0;
-}
-
-// DEBUGGER: r
-// DEBUGGER: p s
-// CHECK: a =
-// DEBUGGER: p s.a[0]
-// CHECK: = 0
-// DEBUGGER: p s.a[1]
-// CHECK: = 1
-// DEBUGGER: p s.a[7]

debuginfo-tests/block_var.m

@@ -1,32 +0,0 @@
-// RUN: %clang %target_itanium_abi_host_triple -O0 -g %s -c -o %t.o
-// RUN: %clang %target_itanium_abi_host_triple %t.o -o %t.out -framework Foundation
-// RUN: %test_debuginfo %s %t.out 
-
-// REQUIRES: system-darwin
-
-// DEBUGGER: break 24
-// DEBUGGER: r
-// DEBUGGER: p result
-// CHECK: ${{[0-9]}} = 42
-
-void doBlock(void (^block)(void))
-{
-    block();
-}
-
-int I(int n)
-{
-    __block int result;
-    int i = 2;
-    doBlock(^{
-        result = n;
-    });
-    return result + i; /* Check value of 'result' */
-}
-
-
-int main (int argc, const char * argv[]) {
-  return I(42);
-}
-
-

debuginfo-tests/blocks.m

@@ -1,43 +0,0 @@
-// RUN: %clang %target_itanium_abi_host_triple -O0 -g %s -c -o %t.o
-// RUN: %clang %target_itanium_abi_host_triple %t.o -o %t.out -framework Foundation
-// RUN: %test_debuginfo %s %t.out 
-
-// REQUIRES: system-darwin
-// Radar 9279956
-
-// DEBUGGER: break 31
-// DEBUGGER: r
-// DEBUGGER: p m2
-// CHECK: ${{[0-9]}} = 1
-// DEBUGGER: p dbTransaction
-// CHECK: ${{[0-9]}} = 0
-// DEBUGGER: p master
-// CHECK: ${{[0-9]}} = 0
-
-#include <Cocoa/Cocoa.h>
-
-extern void foo(void(^)(void));
-
-@interface A:NSObject @end
-@implementation A
-- (void) helper {
- int master = 0;
- __block int m2 = 0;
- __block int dbTransaction = 0;
- int (^x)(void) = ^(void) { (void) self; 
-	(void) master; 
-	(void) dbTransaction; 
-	m2++;
-	return m2;
-	};
-  master = x();
-}
-@end
-
-void foo(void(^x)(void)) {}
-
-int main() {
-	A *a = [A alloc];
-	[a helper];
-	return 0;
-}

debuginfo-tests/ctor.cpp

@@ -1,25 +0,0 @@
-// RUN: %clangxx %target_itanium_abi_host_triple -O0 -g %s -c -o %t.o
-// RUN: %clangxx %target_itanium_abi_host_triple %t.o -o %t.out
-// RUN: %test_debuginfo %s %t.out 
-
-
-// DEBUGGER: break 14
-// DEBUGGER: r
-// DEBUGGER: p *this
-// CHECK-NEXT-NOT: Cannot access memory at address 
-
-class A {
-public:
-	A() : zero(0), data(42)
-	{
-	}
-private:
-	int zero;
-	int data;
-};
-
-int main() {
-	A a;
-	return 0;
-}
-

debuginfo-tests/dbg-arg.c

@@ -1,46 +0,0 @@
-// This test case checks debug info during register moves for an argument.
-// RUN: %clang %target_itanium_abi_host_triple -m64 -mllvm -fast-isel=false  %s -c -o %t.o -g
-// RUN: %clang %target_itanium_abi_host_triple -m64 %t.o -o %t.out
-// RUN: %test_debuginfo %s %t.out
-//
-// DEBUGGER: break 26
-// DEBUGGER: r
-// DEBUGGER: print mutex
-// CHECK:  ={{.* 0x[0-9A-Fa-f]+}}
-//
-// Radar 8412415
-
-struct _mtx
-{
-  long unsigned int ptr;
-  int waiters;
-  struct {
-    int tag;
-    int pad;
-  } mtxi;
-};
-
-
-int foobar(struct _mtx *mutex) {
-  int r = 1;
-  int l = 0;
-  int j = 0;
-  do {
-    if (mutex->waiters) {
-      r = 2;
-    }
-    j = bar(r, l);
-    ++l;
-  } while (l < j);
-  return r + j;
-}
-
-int bar(int i, int j) {
-  return i + j;
-}
-
-int main() {
-  struct _mtx m;
-  m.waiters = 0;
-  return foobar(&m);
-}

debuginfo-tests/foreach.m

@@ -1,31 +0,0 @@
-// RUN: %clang %target_itanium_abi_host_triple -O0 -g %s -c -o %t.o
-// RUN: %clang %target_itanium_abi_host_triple %t.o -o %t.out -framework Foundation
-// RUN: %test_debuginfo %s %t.out 
-//
-// REQUIRES: system-darwin
-// Radar 8757124
-
-// DEBUGGER: break 25
-// DEBUGGER: r
-// DEBUGGER: po thing
-// CHECK: aaa
-
-#import <Foundation/Foundation.h>
-
-int main (int argc, const char * argv[]) {
-
-    NSAutoreleasePool *pool = [[NSAutoreleasePool alloc] init];
-    NSArray *things = [NSArray arrayWithObjects:@"one", @"two", @"three" , nil];
-    for (NSString *thing in things) {
-        NSLog (@"%@", thing);
-    }
-    
-    things = [NSArray arrayWithObjects:@"aaa", @"bbb", @"ccc" , nil];
-    for (NSString *thing in things) {
-        NSLog (@"%@", thing);
-    }
-    [pool release];
-    return 0;
-}
-
-

debuginfo-tests/forward-declare-class.cpp

@@ -1,27 +0,0 @@
-// RUN: %clangxx %target_itanium_abi_host_triple -O0 -g %s -c -o %t.o
-// RUN: %test_debuginfo %s %t.o
-// Radar 9168773
-
-// DEBUGGER: ptype A
-// Work around a gdb bug where it believes that a class is a
-// struct if there aren't any methods - even though it's tagged
-// as a class.
-// CHECK: type = {{struct|class}} A {
-// CHECK-NEXT: {{(public:){0,1}}}
-// CHECK-NEXT: int MyData;
-// CHECK-NEXT: }
-class A;
-class B {
-public:
-  void foo(const A *p);
-};
-
-B iEntry;
-
-class A {
-public:
-  int MyData;
-};
-
-A irp;
-

debuginfo-tests/lit.cfg.py

@@ -1,63 +0,0 @@
-# -*- Python -*-
-
-import os
-import platform
-import re
-import subprocess
-import tempfile
-
-import lit.formats
-import lit.util
-
-from lit.llvm import llvm_config
-from lit.llvm.subst import ToolSubst
-from lit.llvm.subst import FindTool
-
-# Configuration file for the 'lit' test runner.
-
-# name: The name of this test suite.
-config.name = 'debuginfo-tests'
-
-# testFormat: The test format to use to interpret tests.
-#
-# For now we require '&&' between commands, until they get globally killed and
-# the test runner updated.
-config.test_format = lit.formats.ShTest(not llvm_config.use_lit_shell)
-
-# suffixes: A list of file extensions to treat as test files.
-config.suffixes = ['.c', '.cpp', '.m']
-
-# excludes: A list of directories to exclude from the testsuite. The 'Inputs'
-# subdirectories contain auxiliary inputs for various tests in their parent
-# directories.
-config.excludes = ['Inputs']
-
-# test_source_root: The root path where tests are located.
-config.test_source_root = os.path.join(config.debuginfo_tests_src_root)
-
-# test_exec_root: The root path where tests should be run.
-config.test_exec_root = config.debuginfo_tests_obj_root
-
-llvm_config.use_default_substitutions()
-
-# clang_src_dir is not used by these tests, but is required by
-# use_clang(), so set it to "".
-if not hasattr(config, 'clang_src_dir'):
-    config.clang_src_dir = ""
-llvm_config.use_clang()
-
-if config.llvm_use_sanitizer:
-    # Propagate path to symbolizer for ASan/MSan.
-    llvm_config.with_system_environment(
-        ['ASAN_SYMBOLIZER_PATH', 'MSAN_SYMBOLIZER_PATH'])
-
-tool_dirs = [config.llvm_tools_dir]
-
-tools = [
-    ToolSubst('%test_debuginfo', command=os.path.join(
-        config.debuginfo_tests_src_root, 'test_debuginfo.pl')),
-]
-
-llvm_config.add_tool_substitutions(tools, tool_dirs)
-
-lit.util.usePlatformSdkOnDarwin(config, lit_config)

debuginfo-tests/lit.local.cfg

@@ -1,3 +0,0 @@
-# debuginfo-tests are not expected to pass in a cross-compilation setup.
-if 'native' not in config.available_features:
-    config.unsupported = True

debuginfo-tests/lit.site.cfg.py.in

@@ -1,25 +0,0 @@
-@LIT_SITE_CFG_IN_HEADER@
-
-import lit.util
-
-config.test_exec_root = "@CMAKE_BINARY_DIR@"
-
-config.llvm_src_root = "@LLVM_SOURCE_DIR@"
-config.llvm_obj_root = "@LLVM_BINARY_DIR@"
-config.llvm_tools_dir = "@LLVM_TOOLS_DIR@"
-config.llvm_libs_dir = "@LLVM_LIBS_DIR@"
-config.llvm_shlib_dir = "@SHLIBDIR@"
-config.llvm_plugin_ext = "@LLVM_PLUGIN_EXT@"
-config.debuginfo_tests_obj_root = "@DEBUGINFO_TESTS_BINARY_DIR@"
-config.debuginfo_tests_src_root = "@DEBUGINFO_TESTS_SOURCE_DIR@"
-config.has_lld = lit.util.pythonize_bool("@DEBUGINFO_TESTS_HAS_LLD@")
-config.host_triple = "@LLVM_HOST_TRIPLE@"
-config.target_triple = "@TARGET_TRIPLE@"
-config.host_arch = "@HOST_ARCH@"
-
-config.llvm_use_sanitizer = "@LLVM_USE_SANITIZER@"
-
-@LIT_SITE_CFG_IN_FOOTER@
-
-# Let the main config do the real work.
-lit_config.load_config(config, "@DEBUGINFO_TESTS_SOURCE_DIR@/lit.cfg.py")

debuginfo-tests/llgdb.py

@@ -1,157 +0,0 @@
-#!/bin/env python
-"""
-A gdb-compatible frontend for lldb that implements just enough
-commands to run the tests in the debuginfo-tests repository with lldb.
-"""
-
-# ----------------------------------------------------------------------
-# Auto-detect lldb python module.
-import commands, platform, os,  sys
-try:
-    # Just try for LLDB in case PYTHONPATH is already correctly setup.
-    import lldb
-except ImportError:
-    lldb_python_dirs = list()
-    # lldb is not in the PYTHONPATH, try some defaults for the current platform.
-    platform_system = platform.system()
-    if platform_system == 'Darwin':
-        # On Darwin, try the currently selected Xcode directory
-        xcode_dir = commands.getoutput("xcode-select --print-path")
-        if xcode_dir:
-            lldb_python_dirs.append(os.path.realpath(xcode_dir +
-'/../SharedFrameworks/LLDB.framework/Resources/Python'))
-            lldb_python_dirs.append(xcode_dir +
-'/Library/PrivateFrameworks/LLDB.framework/Resources/Python')
-        lldb_python_dirs.append(
-'/System/Library/PrivateFrameworks/LLDB.framework/Resources/Python')
-    success = False
-    for lldb_python_dir in lldb_python_dirs:
-        if os.path.exists(lldb_python_dir):
-            if not (sys.path.__contains__(lldb_python_dir)):
-                sys.path.append(lldb_python_dir)
-                try:
-                    import lldb
-                except ImportError:
-                    pass
-                else:
-                    print 'imported lldb from: "%s"' % (lldb_python_dir)
-                    success = True
-                    break
-    if not success:
-        print "error: couldn't locate the 'lldb' module, please set PYTHONPATH correctly"
-        sys.exit(1)
-# ----------------------------------------------------------------------
-
-# Command line option handling.
-import argparse
-parser = argparse.ArgumentParser(description=__doc__)
-parser.add_argument('--quiet', '-q', action="store_true", help='ignored')
-parser.add_argument('-batch', action="store_true",
-                    help='exit after processing comand line')
-parser.add_argument('-n', action="store_true", help='ignore .lldb file')
-parser.add_argument('-x', dest='script', type=file, help='execute commands from file')
-parser.add_argument("target", help="the program to debug")
-args = parser.parse_args()
-
-
-# Create a new debugger instance.
-debugger = lldb.SBDebugger.Create()
-debugger.SkipLLDBInitFiles(args.n)
-
-# Don't return from lldb function calls until the process stops.
-debugger.SetAsync(False)
-
-# Create a target from a file and arch.
-arch = os.popen("file "+args.target).read().split()[-1]
-target = debugger.CreateTargetWithFileAndArch(args.target, arch)
-
-if not target:
-    print "Could not create target", args.target
-    sys.exit(1)
-
-if not args.script:
-    print "Interactive mode is not implemented."
-    sys.exit(1)
-
-import re
-for command in args.script:
-    # Strip newline and whitespaces and split into words.
-    cmd = command[:-1].strip().split()
-    if not cmd:
-        continue
-
-    print '> %s'% command[:-1]
-
-    try:
-        if re.match('^r|(run)$', cmd[0]):
-            error = lldb.SBError()
-            launchinfo = lldb.SBLaunchInfo([])
-            launchinfo.SetWorkingDirectory(os.getcwd())
-            process = target.Launch(launchinfo, error)
-            print error
-            if not process or error.fail:
-                state = process.GetState()
-                print "State = %d" % state
-                print """
-ERROR: Could not launch process.
-NOTE: There are several reasons why this may happen:
-  * Root needs to run "DevToolsSecurity --enable".
-  * Older versions of lldb cannot launch more than one process simultaneously.
-"""
-                sys.exit(1)
-
-        elif re.match('^b|(break)$', cmd[0]) and len(cmd) == 2:
-            if re.match('[0-9]+', cmd[1]):
-                # b line
-                mainfile = target.FindFunctions('main')[0].compile_unit.file
-                print target.BreakpointCreateByLocation(mainfile, int(cmd[1]))
-            else:
-                # b file:line
-                file, line = cmd[1].split(':')
-                print target.BreakpointCreateByLocation(file, int(line))
-
-        elif re.match('^ptype$', cmd[0]) and len(cmd) == 2:
-            # GDB's ptype has multiple incarnations depending on its
-            # argument (global variable, function, type).  The definition
-            # here is for looking up the signature of a function and only
-            # if that fails it looks for a type with that name.
-            # Type lookup in LLDB would be "image lookup --type".
-            for elem in target.FindFunctions(cmd[1]):
-                print elem.function.type
-                continue
-            print target.FindFirstType(cmd[1])
-
-        elif re.match('^po$', cmd[0]) and len(cmd) > 1:
-            try:
-                opts = lldb.SBExpressionOptions()
-                opts.SetFetchDynamicValue(True)
-                opts.SetCoerceResultToId(True)
-                print target.EvaluateExpression(' '.join(cmd[1:]), opts)
-            except:
-                # FIXME: This is a fallback path for the lab.llvm.org
-                # buildbot running OS X 10.7; it should be removed.
-                thread = process.GetThreadAtIndex(0)
-                frame = thread.GetFrameAtIndex(0)
-                print frame.EvaluateExpression(' '.join(cmd[1:]))
-
-        elif re.match('^p|(print)$', cmd[0]) and len(cmd) > 1:
-            thread = process.GetThreadAtIndex(0)
-            frame = thread.GetFrameAtIndex(0)
-            print frame.EvaluateExpression(' '.join(cmd[1:]))
-
-        elif re.match('^n|(next)$', cmd[0]):
-            thread = process.GetThreadAtIndex(0)
-            thread.StepOver()
-
-        elif re.match('^q|(quit)$', cmd[0]):
-            sys.exit(0)
-
-        else:
-            print debugger.HandleCommand(' '.join(cmd))
-
-    except SystemExit:
-        lldb.SBDebugger_Terminate()
-        raise
-    except:
-        print 'Could not handle the command "%s"' % ' '.join(cmd)
-

debuginfo-tests/nested-struct.cpp

@@ -1,21 +0,0 @@
-// RUN: %clangxx %target_itanium_abi_host_triple -O0 -g %s -c -o %t.o
-// RUN: %test_debuginfo %s %t.o
-// Radar 9440721
-// If debug info for my_number() is emitted outside function foo's scope
-// then a debugger may not be able to handle it. At least one version of
-// gdb crashes in such cases.
-
-// DEBUGGER: ptype foo
-// CHECK: int (void)
-
-int foo() {
-  struct Local {
-    static int my_number() {
-      return 42;
-    }
-  };
-
-  int i = 0;
-  i = Local::my_number();
-  return i + 1;
-}

debuginfo-tests/nrvo-string.cpp

@@ -1,27 +0,0 @@
-// This ensures that DW_OP_deref is inserted when necessary, such as when NRVO
-// of a string object occurs in C++.
-//
-// RUN: %clangxx -O0 -fno-exceptions %target_itanium_abi_host_triple %s -o %t.out -g
-// RUN: %test_debuginfo %s %t.out
-// RUN: %clangxx -O1 -fno-exceptions %target_itanium_abi_host_triple %s -o %t.out -g
-// RUN: %test_debuginfo %s %t.out
-//
-// PR34513
-
-struct string {
-  string() {}
-  string(int i) : i(i) {}
-  ~string() {}
-  int i = 0;
-};
-string get_string() {
-  string unused;
-  string result = 3;
-// DEBUGGER: break 21
-  return result;
-}
-int main() { get_string(); }
-
-// DEBUGGER: r
-// DEBUGGER: print result.i
-// CHECK:  = 3

debuginfo-tests/safestack.c

@@ -1,51 +0,0 @@
-// RUN: %clang %target_itanium_abi_host_triple -arch x86_64 %s -o %t.out -g -fsanitize=safe-stack
-// RUN: %test_debuginfo %s %t.out
-// REQUIRES: not_asan
-//           Zorg configures the ASAN stage2 bots to not build the
-//           safestack compiler-rt.  Only run this test on
-//           non-asanified configurations.
-
-struct S {
-  int a[8];
-};
-
-int f(struct S s, unsigned i);
-
-int main(int argc, const char **argv) {
-  struct S s = {{0, 1, 2, 3, 4, 5, 6, 7}};
-  // DEBUGGER: break 17
-  f(s, 4);
-  // DEBUGGER: break 19
-  return 0;
-}
-
-int f(struct S s, unsigned i) {
-  // DEBUGGER: break 24
-  return s.a[i];
-}
-
-// DEBUGGER: r
-// DEBUGGER: p s
-// CHECK: a =
-// DEBUGGER: p s.a[0]
-// CHECK: = 0
-// DEBUGGER: p s.a[1]
-// CHECK: = 1
-// DEBUGGER: p s.a[7]
-// CHECK: = 7
-// DEBUGGER: c
-// DEBUGGER: p s
-// CHECK: a =
-// DEBUGGER: p s.a[0]
-// CHECK: = 0
-// DEBUGGER: p s.a[1]
-// CHECK: = 1
-// DEBUGGER: p s.a[7]
-// DEBUGGER: c
-// DEBUGGER: p s
-// CHECK: a =
-// DEBUGGER: p s.a[0]
-// CHECK: = 0
-// DEBUGGER: p s.a[1]
-// CHECK: = 1
-// DEBUGGER: p s.a[7]

debuginfo-tests/sret.cpp

@@ -1,71 +0,0 @@
-// RUN: %clangxx %target_itanium_abi_host_triple -O0 -g %s -c -o %t.o
-// RUN: %clangxx %target_itanium_abi_host_triple %t.o -o %t.out
-// RUN: %test_debuginfo %s %t.out 
-// Radar 8775834
-// DEBUGGER: break 62
-// DEBUGGER: r
-// DEBUGGER: p a
-// CHECK: ${{[0-9]+}} =
-// LLDB does not print artificial members.
-// CHECK:  {{(_vptr\$A =)?.*}}m_int = 12
-
-class A
-{
-public:
-    A (int i=0);
-    A (const A& rhs);
-    const A&
-    operator= (const A& rhs);
-    virtual ~A() {}
-
-    int get_int();
-
-protected:
-    int m_int;
-};
-
-A::A (int i) : 
-    m_int(i)
-{
-}
-
-A::A (const A& rhs) :
-    m_int (rhs.m_int)
-{
-}
-
-const A &
-A::operator =(const A& rhs)
-{
-    m_int = rhs.m_int;
-    return *this;
-}
-
-int A::get_int()
-{
-    return m_int;
-}
-
-class B
-{
-public:
-    B () {}
-    
-    A AInstance();
-};
-
-A 
-B::AInstance()
-{
-    A a(12);
-    return a;
-}
-
-int main (int argc, char const *argv[])
-{
-    B b;
-    int return_val = b.AInstance().get_int();
-    
-    A a(b.AInstance());
-    return return_val;
-}

debuginfo-tests/stack-var.c

@@ -1,18 +0,0 @@
-// RUN: %clang %target_itanium_abi_host_triple %s -O -o %t.out -g
-// RUN: %test_debuginfo %s %t.out
-
-void __attribute__((noinline, optnone)) bar(int *test) {}
-int main() {
-  int test;
-  test = 23;
-  // DEBUGGER: break 12
-  // DEBUGGER: r
-  // DEBUGGER: p test
-  // CHECK: = 23
-  bar(&test);
-  // DEBUGGER: break 17
-  // DEBUGGER: c
-  // DEBUGGER: p test
-  // CHECK: = 23
-  return test;
-}

debuginfo-tests/static-member-2.cpp

@@ -1,39 +0,0 @@
-// RUN: %clangxx %target_itanium_abi_host_triple -O0 -g %s -o %t -c
-// RUN: %clangxx %target_itanium_abi_host_triple %t -o %t.out
-// RUN: %test_debuginfo %s %t.out
-
-// FIXME: LLDB finds the wrong symbol for "C". rdar://problem/14933867
-// XFAIL: darwin
-
-// DEBUGGER: delete breakpoints
-// DEBUGGER: break static-member.cpp:33
-// DEBUGGER: r
-// DEBUGGER: ptype C
-// CHECK:      {{struct|class}} C {
-// CHECK:      static const int a;
-// CHECK-NEXT: static int b;
-// CHECK-NEXT: static int c;
-// CHECK-NEXT: int d;
-// CHECK-NEXT: }
-// DEBUGGER: p C::a
-// CHECK: ${{[0-9]}} = 4
-// DEBUGGER: p C::c
-// CHECK: ${{[0-9]}} = 15
-
-// PR14471, PR14734
-
-class C {
-public:
-  const static int a = 4;
-  static int b;
-  static int c;
-  int d;
-};
-
-int C::c = 15;
-const int C::a;
-
-int main() {
-    C instance_C;
-    return C::a;
-}

debuginfo-tests/static-member.cpp

@@ -1,36 +0,0 @@
-// RUN: %clangxx %target_itanium_abi_host_triple -O0 -g %s -o %t -c
-// RUN: %clangxx %target_itanium_abi_host_triple %t -o %t.out
-// RUN: %test_debuginfo %s %t.out
-
-// DEBUGGER: delete breakpoints
-// DEBUGGER: break static-member.cpp:33
-// DEBUGGER: r
-// DEBUGGER: ptype MyClass
-// CHECK:      {{struct|class}} MyClass {
-// CHECK:      static const int a;
-// CHECK-NEXT: static int b;
-// CHECK-NEXT: static int c;
-// CHECK-NEXT: int d;
-// CHECK-NEXT: }
-// DEBUGGER: p MyClass::a
-// CHECK: ${{[0-9]}} = 4
-// DEBUGGER: p MyClass::c
-// CHECK: ${{[0-9]}} = 15
-
-// PR14471, PR14734
-
-class MyClass {
-public:
-  const static int a = 4;
-  static int b;
-  static int c;
-  int d;
-};
-
-int MyClass::c = 15;
-const int MyClass::a;
-
-int main() {
-    MyClass instance_MyClass;
-    return MyClass::a;
-}

debuginfo-tests/test_debuginfo.pl

@@ -1,80 +0,0 @@
-#!/usr/bin/perl
-#
-# This script tests debugging information generated by a compiler.
-# Input arguments
-#   - Input source program. Usually this source file is decorated using
-#     special comments to communicate debugger commands.
-#   - Executable file. This file is generated by the compiler.
-#
-# This perl script extracts debugger commands from input source program
-# comments in a script. A debugger is used to load the executable file
-# and run the script generated from source program comments. Finally,
-# the debugger output is checked, using FileCheck, to validate
-# debugging information.
-#
-# On Darwin the default is to use the llgdb.py wrapper script which
-# translates gdb commands into their lldb equivalents.
-
-use File::Basename;
-use Config;
-use Cwd;
-
-my $testcase_file = $ARGV[0];
-my $executable_file = $ARGV[1];
-
-my $input_filename = basename $testcase_file;
-my $output_dir = dirname $executable_file;
-
-my $debugger_script_file = "$output_dir/$input_filename.debugger.script";
-my $output_file = "$output_dir/$input_filename.gdb.output";
-
-my %cmd_map = ();
-# Assume lldb to be the debugger on Darwin.
-my $use_lldb = 0;
-$use_lldb = 1 if ($Config{osname} eq "darwin");
-
-# Extract debugger commands from testcase. They are marked with DEBUGGER:
-# at the beginning of a comment line.
-open(INPUT, $testcase_file);
-open(OUTPUT, ">$debugger_script_file");
-while(<INPUT>) {
-    my($line) = $_;
-    $i = index($line, "DEBUGGER:");
-    if ( $i >= 0) {
-        $l = length("DEBUGGER:");
-        $s = substr($line, $i + $l);
-        print OUTPUT  "$s";
-    }
-}
-print OUTPUT "\n";
-print OUTPUT "quit\n";
-close(INPUT);
-close(OUTPUT);
-
-# setup debugger and debugger options to run a script.
-my $my_debugger = $ENV{'DEBUGGER'};
-if (!$my_debugger) {
-    if ($use_lldb) {
-        my $path = dirname(Cwd::abs_path($0));
-        $my_debugger = "/usr/bin/env python $path/llgdb.py";
-    } else {
-        $my_debugger = "gdb";
-    }
-}
-
-# quiet / exit after cmdline / no init file / execute script
-my $debugger_options = "-q -batch -n -x";
-
-# run debugger and capture output.
-system("$my_debugger $debugger_options $debugger_script_file $executable_file > $output_file 2>&1");
-
-# validate output.
-system("FileCheck", "-input-file", "$output_file", "$testcase_file");
-if ($?>>8 == 1) {
-    print "Debugger output was:\n";
-    system("cat", "$output_file");
-    exit 1;
-}
-else {
-    exit 0;
-}

debuginfo-tests/vla.c

@@ -1,24 +0,0 @@
-// This test case verifies the debug location for variable-length arrays.
-// RUN: %clang %target_itanium_abi_host_triple -O0 -g %s -c -o %t.o
-// RUN: %clang %target_itanium_abi_host_triple %t.o -o %t.out
-// RUN: %test_debuginfo %s %t.out 
-//
-// DEBUGGER: break 18
-// DEBUGGER: r
-// DEBUGGER: p vla[0]
-// CHECK: 23
-// DEBUGGER: p vla[1]
-// CHECK: 22
-
-void init_vla(int size) {
-  int i;
-  int vla[size];
-  for (i = 0; i < size; i++)
-    vla[i] = size-i;
-  vla[0] = size; // line 18
-}
-
-int main(int argc, const char **argv) {
-  init_vla(23);
-  return 0;
-}

libclc/.gitignore

@@ -1,15 +0,0 @@
-Makefile
-amdgcn--
-amdgcn--amdhsa
-amdgcn-mesa-mesa3d
-build/*.pyc
-built_libs/
-generic--
-generic/lib/convert.cl
-libclc.pc
-nvptx--nvidiacl
-nvptx64--nvidiacl
-r600--
-utils/prepare-builtins
-utils/prepare-builtins.o
-utils/prepare-builtins.o.d

libclc/.travis.yml

@@ -1,64 +0,0 @@
-language: cpp
-
-sudo: false
-dist: trusty
-
-cache:
-  apt: true
-
-
-matrix:
-  include:
-    - env:
-        - LABEL="make gcc LLVM-3.9"
-        - LLVM_VERSION=3.9
-        - LLVM_CONFIG="llvm-config-${LLVM_VERSION}"
-        - CHECK_FILES="barts-r600--.bc cayman-r600--.bc cedar-r600--.bc cypress-r600--.bc tahiti-amdgcn--.bc amdgcn--amdhsa.bc nvptx--nvidiacl.bc nvptx64--nvidiacl.bc"
-      addons:
-        apt:
-          sources:
-            - llvm-toolchain-trusty-3.9
-          packages:
-            - libedit-dev
-            - g++-4.8
-            # From sources above
-            - llvm-3.9-dev
-            - clang-3.9
-    - env:
-        - LABEL="make gcc LLVM-4.0"
-        - LLVM_VERSION=4.0
-        - LLVM_CONFIG="llvm-config-${LLVM_VERSION}"
-        - CHECK_FILES="barts-r600--.bc cayman-r600--.bc cedar-r600--.bc cypress-r600--.bc tahiti-amdgcn--.bc amdgcn--amdhsa.bc nvptx--nvidiacl.bc nvptx64--nvidiacl.bc"
-      addons:
-        apt:
-          sources:
-            - llvm-toolchain-trusty-4.0
-          packages:
-            - libedit-dev
-            - g++-4.8
-            # From sources above
-            - llvm-4.0-dev
-            - clang-4.0
-    - env:
-        - LABEL="make gcc LLVM-5.0"
-        - LLVM_VERSION=5.0
-        - LLVM_CONFIG="llvm-config-${LLVM_VERSION}"
-        - CHECK_FILES="barts-r600--.bc cayman-r600--.bc cedar-r600--.bc cypress-r600--.bc tahiti-amdgcn--.bc amdgcn--amdhsa.bc nvptx--nvidiacl.bc nvptx64--nvidiacl.bc"
-      addons:
-        apt:
-          sources:
-            - llvm-toolchain-trusty-5.0
-          packages:
-            - libedit-dev
-            - g++-4.8
-            # From sources above
-            - llvm-5.0-dev
-            - clang-5.0
-
-script:
-  - $PYTHON ./configure.py --with-llvm-config=$LLVM_CONFIG --with-cxx-compiler=$CXX && make -j4
-  - ret=0;
-    for f in $CHECK_FILES; do
-        ./check_external_calls.sh built_libs/$f || ret=1;
-    done;
-    test $ret -eq 0

libclc/CREDITS.TXT

@@ -1,2 +0,0 @@
-N: Peter Collingbourne
-E: peter@pcc.me.uk

libclc/LICENSE.TXT

@@ -1,64 +0,0 @@
-==============================================================================
-libclc License
-==============================================================================
-
-The libclc library is dual licensed under both the University of Illinois
-"BSD-Like" license and the MIT license.  As a user of this code you may choose
-to use it under either license.  As a contributor, you agree to allow your code
-to be used under both.
-
-Full text of the relevant licenses is included below.
-
-==============================================================================
-
-Copyright (c) 2011-2016 by the contributors listed in CREDITS.TXT
-
-All rights reserved.
-
-Permission is hereby granted, free of charge, to any person obtaining a copy of
-this software and associated documentation files (the "Software"), to deal with
-the Software without restriction, including without limitation the rights to
-use, copy, modify, merge, publish, distribute, sublicense, and/or sell copies
-of the Software, and to permit persons to whom the Software is furnished to do
-so, subject to the following conditions:
-
-    * Redistributions of source code must retain the above copyright notice,
-      this list of conditions and the following disclaimers.
-
-    * Redistributions in binary form must reproduce the above copyright notice,
-      this list of conditions and the following disclaimers in the
-      documentation and/or other materials provided with the distribution.
-
-    * The names of the contributors may not be used to endorse or promote
-      products derived from this Software without specific prior written
-      permission.
-
-THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
-IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, FITNESS
-FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
-CONTRIBUTORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
-LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
-OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS WITH THE
-SOFTWARE.
-
-==============================================================================
-
-Copyright (c) 2011-2014 by the contributors listed in CREDITS.TXT
-
-Permission is hereby granted, free of charge, to any person obtaining a copy
-of this software and associated documentation files (the "Software"), to deal
-in the Software without restriction, including without limitation the rights
-to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
-copies of the Software, and to permit persons to whom the Software is
-furnished to do so, subject to the following conditions:
-
-The above copyright notice and this permission notice shall be included in
-all copies or substantial portions of the Software.
-
-THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
-IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
-FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
-AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
-LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
-OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
-THE SOFTWARE.

libclc/README.TXT

@@ -1,52 +0,0 @@
-libclc
-------
-
-libclc is an open source, BSD licensed implementation of the library
-requirements of the OpenCL C programming language, as specified by the
-OpenCL 1.1 Specification. The following sections of the specification
-impose library requirements:
-
-  * 6.1: Supported Data Types
-  * 6.2.3: Explicit Conversions
-  * 6.2.4.2: Reinterpreting Types Using as_type() and as_typen()
-  * 6.9: Preprocessor Directives and Macros
-  * 6.11: Built-in Functions
-  * 9.3: Double Precision Floating-Point
-  * 9.4: 64-bit Atomics
-  * 9.5: Writing to 3D image memory objects
-  * 9.6: Half Precision Floating-Point
-
-libclc is intended to be used with the Clang compiler's OpenCL frontend.
-
-libclc is designed to be portable and extensible. To this end, it provides
-generic implementations of most library requirements, allowing the target
-to override the generic implementation at the granularity of individual
-functions.
-
-libclc currently only supports the PTX target, but support for more
-targets is welcome.
-
-Compiling and installing with Make
-----------------------------------
-
-$ ./configure.py --with-llvm-config=/path/to/llvm-config && make
-$ make install
-
-Note you can use the DESTDIR Makefile variable to do staged installs.
-
-$ make install DESTDIR=/path/for/staged/install
-
-Compiling and installing with Ninja
------------------------------------
-
-$ ./configure.py -g ninja --with-llvm-config=/path/to/llvm-config && ninja
-$ ninja install
-
-Note you can use the DESTDIR environment variable to do staged installs.
-
-$ DESTDIR=/path/for/staged/install ninja install
-
-Website
--------
-
-http://www.pcc.me.uk/~peter/libclc/

libclc/amdgcn-amdhsa/lib/OVERRIDES

@@ -1 +0,0 @@
-workitem/get_num_groups.ll

libclc/amdgcn-amdhsa/lib/OVERRIDES_3.9

@@ -1,3 +0,0 @@
-workitem/get_global_size.ll
-workitem/get_local_size.ll
-workitem/get_num_groups.39.ll

libclc/amdgcn-amdhsa/lib/SOURCES

@@ -1,3 +0,0 @@
-workitem/get_global_size.ll
-workitem/get_local_size.ll
-workitem/get_num_groups.cl

libclc/amdgcn-amdhsa/lib/SOURCES_3.9

@@ -1,2 +0,0 @@
-workitem/get_global_size.39.ll
-workitem/get_local_size.39.ll

libclc/amdgcn-amdhsa/lib/workitem/get_global_size.39.ll

@@ -1,36 +0,0 @@
-declare i8 addrspace(2)* @llvm.amdgcn.dispatch.ptr() #0
-
-define i32 @get_global_size(i32 %dim) #1 {
-  %dispatch_ptr = call noalias nonnull dereferenceable(64) i8 addrspace(2)* @llvm.amdgcn.dispatch.ptr()
-  switch i32 %dim, label %default [
-    i32 0, label %x
-    i32 1, label %y
-    i32 2, label %z
-  ]
-
-x:
-  %ptr_x = getelementptr inbounds i8, i8 addrspace(2)* %dispatch_ptr, i32 12
-  %ptr_x32 = bitcast i8 addrspace(2)* %ptr_x to i32 addrspace(2)*
-  %x32 = load i32, i32 addrspace(2)* %ptr_x32, align 4, !invariant.load !0
-  ret i32 %x32
-
-y:
-  %ptr_y = getelementptr inbounds i8, i8 addrspace(2)* %dispatch_ptr, i32 16
-  %ptr_y32 = bitcast i8 addrspace(2)* %ptr_y to i32 addrspace(2)*
-  %y32 = load i32, i32 addrspace(2)* %ptr_y32, align 4, !invariant.load !0
-  ret i32 %y32
-
-z:
-  %ptr_z = getelementptr inbounds i8, i8 addrspace(2)* %dispatch_ptr, i32 20
-  %ptr_z32 = bitcast i8 addrspace(2)* %ptr_z to i32 addrspace(2)*
-  %z32 = load i32, i32 addrspace(2)* %ptr_z32, align 4, !invariant.load !0
-  ret i32 %z32
-
-default:
-  ret i32 1
-}
-
-attributes #0 = { nounwind readnone }
-attributes #1 = { alwaysinline norecurse nounwind readonly }
-
-!0 = !{}

libclc/amdgcn-amdhsa/lib/workitem/get_global_size.ll

@@ -1,39 +0,0 @@
-declare i8 addrspace(2)* @llvm.amdgcn.dispatch.ptr() #0
-
-define i64 @get_global_size(i32 %dim) #1 {
-  %dispatch_ptr = call noalias nonnull dereferenceable(64) i8 addrspace(2)* @llvm.amdgcn.dispatch.ptr()
-  switch i32 %dim, label %default [
-    i32 0, label %x
-    i32 1, label %y
-    i32 2, label %z
-  ]
-
-x:
-  %ptr_x = getelementptr inbounds i8, i8 addrspace(2)* %dispatch_ptr, i64 12
-  %ptr_x32 = bitcast i8 addrspace(2)* %ptr_x to i32 addrspace(2)*
-  %x32 = load i32, i32 addrspace(2)* %ptr_x32, align 4, !invariant.load !0
-  %size_x = zext i32 %x32 to i64
-  ret i64 %size_x
-
-y:
-  %ptr_y = getelementptr inbounds i8, i8 addrspace(2)* %dispatch_ptr, i64 16
-  %ptr_y32 = bitcast i8 addrspace(2)* %ptr_y to i32 addrspace(2)*
-  %y32 = load i32, i32 addrspace(2)* %ptr_y32, align 4, !invariant.load !0
-  %size_y = zext i32 %y32 to i64
-  ret i64 %size_y
-
-z:
-  %ptr_z = getelementptr inbounds i8, i8 addrspace(2)* %dispatch_ptr, i64 20
-  %ptr_z32 = bitcast i8 addrspace(2)* %ptr_z to i32 addrspace(2)*
-  %z32 = load i32, i32 addrspace(2)* %ptr_z32, align 4, !invariant.load !0
-  %size_z = zext i32 %z32 to i64
-  ret i64 %size_z
-
-default:
-  ret i64 1
-}
-
-attributes #0 = { nounwind readnone }
-attributes #1 = { alwaysinline norecurse nounwind readonly }
-
-!0 = !{}

libclc/amdgcn-amdhsa/lib/workitem/get_local_size.39.ll

@@ -1,35 +0,0 @@
-declare i8 addrspace(2)* @llvm.amdgcn.dispatch.ptr() #0
-
-define i32 @get_local_size(i32 %dim) #1 {
-  %dispatch_ptr = call noalias nonnull dereferenceable(64) i8 addrspace(2)* @llvm.amdgcn.dispatch.ptr()
-  %dispatch_ptr_i32 = bitcast i8 addrspace(2)* %dispatch_ptr to i32 addrspace(2)*
-  %xy_size_ptr = getelementptr inbounds i32, i32 addrspace(2)* %dispatch_ptr_i32, i32 1
-  %xy_size = load i32, i32 addrspace(2)* %xy_size_ptr, align 4, !invariant.load !0
-  switch i32 %dim, label %default [
-    i32 0, label %x_dim
-    i32 1, label %y_dim
-    i32 2, label %z_dim
-  ]
-
-x_dim:
-  %x_size = and i32 %xy_size, 65535
-  ret i32 %x_size
-
-y_dim:
-  %y_size = lshr i32 %xy_size, 16
-  ret i32 %y_size
-
-z_dim:
-  %z_size_ptr = getelementptr inbounds i32, i32 addrspace(2)* %dispatch_ptr_i32, i32 2
-  %z_size = load i32, i32 addrspace(2)* %z_size_ptr, align 4, !invariant.load !0, !range !1
-  ret i32 %z_size
-
-default:
-  ret i32 1
-}
-
-attributes #0 = { nounwind readnone }
-attributes #1 = { alwaysinline norecurse nounwind readonly }
-
-!0 = !{}
-!1 = !{ i32 0, i32 257 }

libclc/amdgcn-amdhsa/lib/workitem/get_local_size.ll

@@ -1,38 +0,0 @@
-declare i8 addrspace(2)* @llvm.amdgcn.dispatch.ptr() #0
-
-define i64 @get_local_size(i32 %dim) #1 {
-  %dispatch_ptr = call noalias nonnull dereferenceable(64) i8 addrspace(2)* @llvm.amdgcn.dispatch.ptr()
-  %dispatch_ptr_i32 = bitcast i8 addrspace(2)* %dispatch_ptr to i32 addrspace(2)*
-  %xy_size_ptr = getelementptr inbounds i32, i32 addrspace(2)* %dispatch_ptr_i32, i64 1
-  %xy_size = load i32, i32 addrspace(2)* %xy_size_ptr, align 4, !invariant.load !0
-  switch i32 %dim, label %default [
-    i32 0, label %x_dim
-    i32 1, label %y_dim
-    i32 2, label %z_dim
-  ]
-
-x_dim:
-  %x_size = and i32 %xy_size, 65535
-  %x_size.ext = zext i32 %x_size to i64
-  ret i64 %x_size.ext
-
-y_dim:
-  %y_size = lshr i32 %xy_size, 16
-  %y_size.ext = zext i32 %y_size to i64
-  ret i64 %y_size.ext
-
-z_dim:
-  %z_size_ptr = getelementptr inbounds i32, i32 addrspace(2)* %dispatch_ptr_i32, i64 2
-  %z_size = load i32, i32 addrspace(2)* %z_size_ptr, align 4, !invariant.load !0, !range !1
-  %z_size.ext = zext i32 %z_size to i64
-  ret i64 %z_size.ext
-
-default:
-  ret i64 1
-}
-
-attributes #0 = { nounwind readnone }
-attributes #1 = { alwaysinline norecurse nounwind readonly }
-
-!0 = !{}
-!1 = !{ i32 0, i32 257 }

libclc/amdgcn-amdhsa/lib/workitem/get_num_groups.cl

@@ -1,12 +0,0 @@
-
-#include <clc/clc.h>
-
-_CLC_DEF size_t get_num_groups(uint dim) {
-  size_t global_size = get_global_size(dim);
-  size_t local_size = get_local_size(dim);
-  size_t num_groups = global_size / local_size;
-  if (global_size % local_size != 0) {
-    num_groups++;
-  }
-  return num_groups;
-}

libclc/amdgcn/lib/OVERRIDES_3.9

@@ -1,3 +0,0 @@
-workitem/get_global_size.ll
-workitem/get_local_size.ll
-workitem/get_num_groups.ll

libclc/amdgcn/lib/SOURCES

@@ -1,12 +0,0 @@
-cl_khr_int64_extended_atomics/minmax_helpers.ll
-math/ldexp.cl
-mem_fence/fence.cl
-mem_fence/waitcnt.ll
-synchronization/barrier.cl
-workitem/get_global_offset.cl
-workitem/get_group_id.cl
-workitem/get_global_size.ll
-workitem/get_local_id.cl
-workitem/get_local_size.ll
-workitem/get_num_groups.ll
-workitem/get_work_dim.cl

libclc/amdgcn/lib/SOURCES_3.9

@@ -1,3 +0,0 @@
-workitem/get_global_size.39.ll
-workitem/get_local_size.39.ll
-workitem/get_num_groups.39.ll

libclc/amdgcn/lib/cl_khr_int64_extended_atomics/minmax_helpers.ll

@@ -1,49 +0,0 @@
-target datalayout = "e-p:32:32-p1:64:64-p2:64:64-p3:32:32-p4:64:64-p5:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-v2048:2048-n32:64"
-
-define i64 @__clc__sync_fetch_and_min_global_8(i64 addrspace(1)* nocapture %ptr, i64 %value) nounwind alwaysinline {
-entry:
-  %0 = atomicrmw volatile min i64 addrspace(1)* %ptr, i64 %value seq_cst
-  ret i64 %0
-}
-
-define i64 @__clc__sync_fetch_and_umin_global_8(i64 addrspace(1)* nocapture %ptr, i64 %value) nounwind alwaysinline {
-entry:
-  %0 = atomicrmw volatile umin i64 addrspace(1)* %ptr, i64 %value seq_cst
-  ret i64 %0
-}
-
-define i64 @__clc__sync_fetch_and_min_local_8(i64 addrspace(3)* nocapture %ptr, i64 %value) nounwind alwaysinline {
-entry:
-  %0 = atomicrmw volatile min i64 addrspace(3)* %ptr, i64 %value seq_cst
-  ret i64 %0
-}
-
-define i64 @__clc__sync_fetch_and_umin_local_8(i64 addrspace(3)* nocapture %ptr, i64 %value) nounwind alwaysinline {
-entry:
-  %0 = atomicrmw volatile umin i64 addrspace(3)* %ptr, i64 %value seq_cst
-  ret i64 %0
-}
-
-define i64 @__clc__sync_fetch_and_max_global_8(i64 addrspace(1)* nocapture %ptr, i64 %value) nounwind alwaysinline {
-entry:
-  %0 = atomicrmw volatile max i64 addrspace(1)* %ptr, i64 %value seq_cst
-  ret i64 %0
-}
-
-define i64 @__clc__sync_fetch_and_umax_global_8(i64 addrspace(1)* nocapture %ptr, i64 %value) nounwind alwaysinline {
-entry:
-  %0 = atomicrmw volatile umax i64 addrspace(1)* %ptr, i64 %value seq_cst
-  ret i64 %0
-}
-
-define i64 @__clc__sync_fetch_and_max_local_8(i64 addrspace(3)* nocapture %ptr, i64 %value) nounwind alwaysinline {
-entry:
-  %0 = atomicrmw volatile max i64 addrspace(3)* %ptr, i64 %value seq_cst
-  ret i64 %0
-}
-
-define i64 @__clc__sync_fetch_and_umax_local_8(i64 addrspace(3)* nocapture %ptr, i64 %value) nounwind alwaysinline {
-entry:
-  %0 = atomicrmw volatile umax i64 addrspace(3)* %ptr, i64 %value seq_cst
-  ret i64 %0
-}

libclc/amdgcn/lib/math/ldexp.cl

@@ -1,47 +0,0 @@
-/*
- * Copyright (c) 2014 Advanced Micro Devices, Inc.
- *
- * Permission is hereby granted, free of charge, to any person obtaining a copy
- * of this software and associated documentation files (the "Software"), to deal
- * in the Software without restriction, including without limitation the rights
- * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
- * copies of the Software, and to permit persons to whom the Software is
- * furnished to do so, subject to the following conditions:
- *
- * The above copyright notice and this permission notice shall be included in
- * all copies or substantial portions of the Software.
- *
- * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
- * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
- * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
- * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
- * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
- * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
- * THE SOFTWARE.
- */
-
-#include <clc/clc.h>
-
-#include "../../../generic/lib/clcmacro.h"
-
-#ifdef __HAS_LDEXPF__
-#define BUILTINF __builtin_amdgcn_ldexpf
-#else
-#include "math/clc_ldexp.h"
-#define BUILTINF __clc_ldexp
-#endif
-
-// This defines all the ldexp(floatN, intN) variants.
-_CLC_DEFINE_BINARY_BUILTIN(float, ldexp, BUILTINF, float, int);
-
-#ifdef cl_khr_fp64
-  #pragma OPENCL EXTENSION cl_khr_fp64 : enable
-    // This defines all the ldexp(doubleN, intN) variants.
-  _CLC_DEFINE_BINARY_BUILTIN(double, ldexp, __builtin_amdgcn_ldexp, double, int);
-#endif
-
-// This defines all the ldexp(GENTYPE, int);
-#define __CLC_BODY <../../../generic/lib/math/ldexp.inc>
-#include <clc/math/gentype.inc>
-
-#undef BUILTINF

libclc/amdgcn/lib/mem_fence/fence.cl

@@ -1,39 +0,0 @@
-#include <clc/clc.h>
-
-void __clc_amdgcn_s_waitcnt(unsigned flags);
-
-// s_waitcnt takes 16bit argument with a combined number of maximum allowed
-// pending operations:
-// [12:8] LGKM -- LDS, GDS, Konstant (SMRD), Messages
-// [7] -- undefined
-// [6:4] -- exports, GDS, and mem write
-// [3:0] -- vector memory operations
-
-// Newer clang supports __builtin_amdgcn_s_waitcnt
-#if __clang_major__ >= 5
-#  define __waitcnt(x) __builtin_amdgcn_s_waitcnt(x)
-#else
-#  define __waitcnt(x) __clc_amdgcn_s_waitcnt(x)
-#endif
-
-_CLC_DEF void mem_fence(cl_mem_fence_flags flags)
-{
-	if (flags & CLK_GLOBAL_MEM_FENCE) {
-		// scalar loads are counted with LGKM but we don't know whether
-		// the compiler turned any loads to scalar
-		__waitcnt(0);
-	} else if (flags & CLK_LOCAL_MEM_FENCE)
-		__waitcnt(0xff); // LGKM is [12:8]
-}
-#undef __waitcnt
-
-// We don't have separate mechanism for read and write fences
-_CLC_DEF void read_mem_fence(cl_mem_fence_flags flags)
-{
-	mem_fence(flags);
-}
-
-_CLC_DEF void write_mem_fence(cl_mem_fence_flags flags)
-{
-	mem_fence(flags);
-}

libclc/amdgcn/lib/mem_fence/waitcnt.ll

@@ -1,13 +0,0 @@
-declare void @llvm.amdgcn.s.waitcnt(i32) #0
-
-target datalayout = "e-p:32:32-p1:64:64-p2:64:64-p3:32:32-p4:64:64-p5:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-v2048:2048-n32:64"
-
-; Export waitcnt intrinsic for clang < 5
-define void @__clc_amdgcn_s_waitcnt(i32 %flags) #1 {
-entry:
-  tail call void @llvm.amdgcn.s.waitcnt(i32 %flags)
-  ret void
-}
-
-attributes #0 = { nounwind }
-attributes #1 = { nounwind alwaysinline }

libclc/amdgcn/lib/synchronization/barrier.cl

@@ -1,7 +0,0 @@
-#include <clc/clc.h>
-
-_CLC_DEF void barrier(cl_mem_fence_flags flags)
-{
-	mem_fence(flags);
-	__builtin_amdgcn_s_barrier();
-}

libclc/amdgcn/lib/workitem/get_global_offset.cl

@@ -1,11 +0,0 @@
-#include <clc/clc.h>
-
-_CLC_DEF size_t get_global_offset(uint dim)
-{
-	__attribute__((address_space(2))) uint * ptr =
-		(__attribute__((address_space(2))) uint *)
-		__builtin_amdgcn_implicitarg_ptr();
-	if (dim < 3)
-		return ptr[dim + 1];
-	return 0;
-}

libclc/amdgcn/lib/workitem/get_global_size.39.ll

@@ -1,20 +0,0 @@
-declare i32 @llvm.r600.read.global.size.x() nounwind readnone
-declare i32 @llvm.r600.read.global.size.y() nounwind readnone
-declare i32 @llvm.r600.read.global.size.z() nounwind readnone
-
-target datalayout = "e-p:32:32-p1:64:64-p2:64:64-p3:32:32-p4:64:64-p5:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-v2048:2048-n32:64"
-
-define i32 @get_global_size(i32 %dim) nounwind readnone alwaysinline {
-  switch i32 %dim, label %default [i32 0, label %x_dim i32 1, label %y_dim i32 2, label %z_dim]
-x_dim:
-  %x = call i32 @llvm.r600.read.global.size.x()
-  ret i32 %x
-y_dim:
-  %y = call i32 @llvm.r600.read.global.size.y()
-  ret i32 %y
-z_dim:
-  %z = call i32 @llvm.r600.read.global.size.z()
-  ret i32 %z
-default:
-  ret i32 1
-}

libclc/amdgcn/lib/workitem/get_global_size.ll

@@ -1,23 +0,0 @@
-declare i32 @llvm.r600.read.global.size.x() nounwind readnone
-declare i32 @llvm.r600.read.global.size.y() nounwind readnone
-declare i32 @llvm.r600.read.global.size.z() nounwind readnone
-
-target datalayout = "e-p:32:32-p1:64:64-p2:64:64-p3:32:32-p4:64:64-p5:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-v2048:2048-n32:64"
-
-define i64 @get_global_size(i32 %dim) nounwind readnone alwaysinline {
-  switch i32 %dim, label %default [i32 0, label %x_dim i32 1, label %y_dim i32 2, label %z_dim]
-x_dim:
-  %x = call i32 @llvm.r600.read.global.size.x()
-  %x.ext = zext i32 %x to i64
-  ret i64 %x.ext
-y_dim:
-  %y = call i32 @llvm.r600.read.global.size.y()
-  %y.ext = zext i32 %y to i64
-  ret i64 %y.ext
-z_dim:
-  %z = call i32 @llvm.r600.read.global.size.z()
-  %z.ext = zext i32 %z to i64
-  ret i64 %z.ext
-default:
-  ret i64 1
-}

libclc/amdgcn/lib/workitem/get_group_id.cl

@@ -1,11 +0,0 @@
-#include <clc/clc.h>
-
-_CLC_DEF size_t get_group_id(uint dim)
-{
-	switch(dim) {
-	case 0: return __builtin_amdgcn_workgroup_id_x();
-	case 1: return __builtin_amdgcn_workgroup_id_y();
-	case 2: return __builtin_amdgcn_workgroup_id_z();
-	default: return 1;
-	}
-}

libclc/amdgcn/lib/workitem/get_local_id.cl

@@ -1,11 +0,0 @@
-#include <clc/clc.h>
-
-_CLC_DEF size_t get_local_id(uint dim)
-{
-	switch(dim) {
-	case 0: return __builtin_amdgcn_workitem_id_x();
-	case 1: return __builtin_amdgcn_workitem_id_y();
-	case 2: return __builtin_amdgcn_workitem_id_z();
-	default: return 1;
-	}
-}

libclc/amdgcn/lib/workitem/get_local_size.39.ll

@@ -1,20 +0,0 @@
-declare i32 @llvm.r600.read.local.size.x() nounwind readnone
-declare i32 @llvm.r600.read.local.size.y() nounwind readnone
-declare i32 @llvm.r600.read.local.size.z() nounwind readnone
-
-target datalayout = "e-p:32:32-p1:64:64-p2:64:64-p3:32:32-p4:64:64-p5:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-v2048:2048-n32:64"
-
-define i32 @get_local_size(i32 %dim) nounwind readnone alwaysinline {
-  switch i32 %dim, label %default [i32 0, label %x_dim i32 1, label %y_dim i32 2, label %z_dim]
-x_dim:
-  %x = call i32 @llvm.r600.read.local.size.x()
-  ret i32 %x
-y_dim:
-  %y = call i32 @llvm.r600.read.local.size.y()
-  ret i32 %y
-z_dim:
-  %z = call i32 @llvm.r600.read.local.size.z()
-  ret i32 %z
-default:
-  ret i32 1
-}

libclc/amdgcn/lib/workitem/get_local_size.ll

@@ -1,23 +0,0 @@
-declare i32 @llvm.r600.read.local.size.x() nounwind readnone
-declare i32 @llvm.r600.read.local.size.y() nounwind readnone
-declare i32 @llvm.r600.read.local.size.z() nounwind readnone
-
-target datalayout = "e-p:32:32-p1:64:64-p2:64:64-p3:32:32-p4:64:64-p5:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-v2048:2048-n32:64"
-
-define i64 @get_local_size(i32 %dim) nounwind readnone alwaysinline {
-  switch i32 %dim, label %default [i32 0, label %x_dim i32 1, label %y_dim i32 2, label %z_dim]
-x_dim:
-  %x = call i32 @llvm.r600.read.local.size.x()
-  %x.ext = zext i32 %x to i64
-  ret i64 %x.ext
-y_dim:
-  %y = call i32 @llvm.r600.read.local.size.y()
-  %y.ext = zext i32 %y to i64
-  ret i64 %y.ext
-z_dim:
-  %z = call i32 @llvm.r600.read.local.size.z()
-  %z.ext = zext i32 %z to i64
-  ret i64 %z.ext
-default:
-  ret i64 1
-}

libclc/amdgcn/lib/workitem/get_num_groups.39.ll

@@ -1,20 +0,0 @@
-declare i32 @llvm.r600.read.ngroups.x() nounwind readnone
-declare i32 @llvm.r600.read.ngroups.y() nounwind readnone
-declare i32 @llvm.r600.read.ngroups.z() nounwind readnone
-
-target datalayout = "e-p:32:32-p1:64:64-p2:64:64-p3:32:32-p4:64:64-p5:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-v2048:2048-n32:64"
-
-define i32 @get_num_groups(i32 %dim) nounwind readnone alwaysinline {
-  switch i32 %dim, label %default [i32 0, label %x_dim i32 1, label %y_dim i32 2, label %z_dim]
-x_dim:
-  %x = call i32 @llvm.r600.read.ngroups.x()
-  ret i32 %x
-y_dim:
-  %y = call i32 @llvm.r600.read.ngroups.y()
-  ret i32 %y
-z_dim:
-  %z = call i32 @llvm.r600.read.ngroups.z()
-  ret i32 %z
-default:
-  ret i32 1
-}

libclc/amdgcn/lib/workitem/get_num_groups.ll

@@ -1,23 +0,0 @@
-declare i32 @llvm.r600.read.ngroups.x() nounwind readnone
-declare i32 @llvm.r600.read.ngroups.y() nounwind readnone
-declare i32 @llvm.r600.read.ngroups.z() nounwind readnone
-
-target datalayout = "e-p:32:32-p1:64:64-p2:64:64-p3:32:32-p4:64:64-p5:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-v2048:2048-n32:64"
-
-define i64 @get_num_groups(i32 %dim) nounwind readnone alwaysinline {
-  switch i32 %dim, label %default [i32 0, label %x_dim i32 1, label %y_dim i32 2, label %z_dim]
-x_dim:
-  %x = call i32 @llvm.r600.read.ngroups.x()
-  %x.ext = zext i32 %x to i64
-  ret i64 %x.ext
-y_dim:
-  %y = call i32 @llvm.r600.read.ngroups.y()
-  %y.ext = zext i32 %y to i64
-  ret i64 %y.ext
-z_dim:
-  %z = call i32 @llvm.r600.read.ngroups.z()
-  %z.ext = zext i32 %z to i64
-  ret i64 %z.ext
-default:
-  ret i64 1
-}

libclc/amdgcn/lib/workitem/get_work_dim.cl

@@ -1,9 +0,0 @@
-#include <clc/clc.h>
-
-_CLC_DEF uint get_work_dim(void)
-{
-	__attribute__((address_space(2))) uint * ptr =
-		(__attribute__((address_space(2))) uint *)
-		__builtin_amdgcn_implicitarg_ptr();
-	return ptr[0];
-}

libclc/amdgpu/lib/OVERRIDES

@@ -1,2 +0,0 @@
-workitem/get_group_id.cl
-workitem/get_global_size.cl

libclc/amdgpu/lib/SOURCES

@@ -1,5 +0,0 @@
-math/native_exp.cl
-math/native_log.cl
-math/native_log10.cl
-math/nextafter.cl
-math/sqrt.cl