Merging r308978:

------------------------------------------------------------------------
r308978 | fedor.sergeev | 2017-07-25 08:28:28 -0700 (Tue, 25 Jul 2017) | 16 lines

[Sparc] invalid adjustments in TLS_LE/TLS_LDO relocations removed

Summary:
Some SPARC TLS relocations were applying nontrivial adjustments
to zero value, leading to unexpected non-zero values in ELF and then
Solaris linker failures.

Getting rid of these adjustments.

Fixes PR33825.

Reviewers: rafael, asb, jyknight

Subscribers: joerg, jyknight, llvm-commits

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

llvm-svn: 309187

clang/docs/ClangCommandLineReference.rst

@@ -96,6 +96,8 @@ Emit ARC errors even if the migrator can fix them
 
 Output path for the plist report
 
+.. option:: --autocomplete=<arg>
+
 .. option:: -bind\_at\_load
 
 .. option:: -bundle
@@ -292,7 +294,7 @@ Disable builtin #include directories
 
 .. option:: -nomultidefs
 
-.. option:: -nopie
+.. option:: -nopie, -no-pie
 
 .. option:: -noprebind
 
@@ -704,6 +706,10 @@ Don't use blacklist file for sanitizers
 
 Level of field padding for AddressSanitizer
 
+.. option:: -fsanitize-address-globals-dead-stripping
+
+Enable linker dead stripping of globals in AddressSanitizer
+
 .. option:: -fsanitize-address-use-after-scope, -fno-sanitize-address-use-after-scope
 
 Enable use-after-scope detection in AddressSanitizer
@@ -1071,6 +1077,10 @@ Target-independent compilation options
 
 Enable C++17 aligned allocation functions
 
+.. option:: -fallow-editor-placeholders, -fno-allow-editor-placeholders
+
+Treat editor placeholders as valid source code
+
 .. option:: -fallow-unsupported
 
 .. option:: -faltivec, -fno-altivec
@@ -1205,6 +1215,10 @@ Print absolute paths in diagnostics
 .. option:: -fdiagnostics-color=<arg>
 .. program:: clang
 
+.. option:: -fdiagnostics-hotness-threshold=<number>
+
+Prevent optimization remarks from being output if they do not have at least this profile count
+
 .. option:: -fdiagnostics-show-hotness, -fno-diagnostics-show-hotness
 
 Enable profile hotness information in diagnostic line
@@ -1585,6 +1599,8 @@ Turn on loop reroller
 
 .. option:: -fretain-comments-from-system-headers
 
+.. option:: -frewrite-imports, -fno-rewrite-imports
+
 .. option:: -frewrite-includes, -fno-rewrite-includes
 
 .. option:: -frewrite-map-file <arg>
@@ -1639,10 +1655,6 @@ Use SjLj style exceptions
 
 Enable the superword-level parallelism vectorization passes
 
-.. option:: -fslp-vectorize-aggressive, -fno-slp-vectorize-aggressive
-
-Enable the BB vectorization passes
-
 .. option:: -fspell-checking, -fno-spell-checking
 
 .. option:: -fspell-checking-limit=<arg>
@@ -1911,6 +1923,8 @@ Link stack frames through backchain on System Z
 
 .. 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:: -mdefault-build-attributes<arg>, -mno-default-build-attributes<arg>
+
 .. option:: -mdll<arg>
 
 .. option:: -mdouble-float
@@ -1975,10 +1989,16 @@ Use Intel MCU ABI
 
 Generate branches with extended addressability, usually via indirect jumps.
 
-.. option:: -mmacosx-version-min=<arg>
+.. option:: -mmacosx-version-min=<arg>, -mmacos-version-min=<arg>
 
 Set Mac OS X deployment target
 
+.. option:: -mmadd4, -mno-madd4
+
+Enable the generation of 4-operand madd.s, madd.d and related instructions.
+
+.. option:: -mmcu=<arg>
+
 .. option:: -mmicromips, -mno-micromips
 
 .. option:: -mms-bitfields, -mno-ms-bitfields
@@ -1989,6 +2009,10 @@ Set the default structure layout to be compatible with the Microsoft compiler st
 
 Enable MSA ASE (MIPS only)
 
+.. option:: -mmt, -mno-mt
+
+Enable MT ASE (MIPS only)
+
 .. option:: -mnan=<arg>
 
 .. option:: -mno-mips16
@@ -2203,6 +2227,8 @@ X86
 
 .. option:: -mavx512vl, -mno-avx512vl
 
+.. option:: -mavx512vpopcntdq, -mno-avx512vpopcntdq
+
 .. option:: -mbmi, -mno-bmi
 
 .. option:: -mbmi2, -mno-bmi2
@@ -2225,6 +2251,8 @@ X86
 
 .. option:: -mfxsr, -mno-fxsr
 
+.. option:: -mlwp, -mno-lwp
+
 .. option:: -mlzcnt, -mno-lzcnt
 
 .. option:: -mmmx, -mno-mmx
@@ -2372,6 +2400,16 @@ Debug information flags
 
 .. option:: -gstrict-dwarf, -gno-strict-dwarf
 
+.. option:: -gz
+
+DWARF debug sections compression type
+
+.. program:: clang1
+.. option:: -gz=<arg>
+.. program:: clang
+
+DWARF debug sections compression type
+
 Static analyzer flags
 =====================
 

clang/include/clang/AST/Type.h

@@ -1752,6 +1752,7 @@ public:
   bool isTemplateTypeParmType() const;          // C++ template type parameter
   bool isNullPtrType() const;                   // C++11 std::nullptr_t
   bool isAlignValT() const;                     // C++17 std::align_val_t
+  bool isStdByteType() const;                   // C++17 std::byte
   bool isAtomicType() const;                    // C11 _Atomic()
 
 #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \

clang/include/clang/Basic/DiagnosticGroups.td

@@ -469,9 +469,8 @@ def IgnoredPragmaIntrinsic : DiagGroup<"ignored-pragma-intrinsic">;
 def UnknownPragmas : DiagGroup<"unknown-pragmas">;
 def IgnoredPragmas : DiagGroup<"ignored-pragmas", [IgnoredPragmaIntrinsic]>;
 def PragmaClangAttribute : DiagGroup<"pragma-clang-attribute">;
-def PragmaPack : DiagGroup<"pragma-pack">;
 def Pragmas : DiagGroup<"pragmas", [UnknownPragmas, IgnoredPragmas,
-                                    PragmaClangAttribute, PragmaPack]>;
+                                    PragmaClangAttribute]>;
 def UnknownWarningOption : DiagGroup<"unknown-warning-option">;
 def NSobjectAttribute : DiagGroup<"NSObject-attribute">;
 def IndependentClassAttribute : DiagGroup<"IndependentClass-attribute">;

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -712,16 +712,6 @@ def err_pragma_options_align_mac68k_target_unsupported : Error<
 def warn_pragma_pack_invalid_alignment : Warning<
   "expected #pragma pack parameter to be '1', '2', '4', '8', or '16'">,
   InGroup<IgnoredPragmas>;
-def warn_pragma_pack_non_default_at_include : Warning<
-  "non-default #pragma pack value might change the alignment of struct or "
-  "union members in the included file">, InGroup<PragmaPack>;
-def warn_pragma_pack_modified_after_include : Warning<
-  "the current #pragma pack aligment value is modified in the included "
-  "file">, InGroup<PragmaPack>;
-def warn_pragma_pack_no_pop_eof : Warning<"unterminated "
-  "'#pragma pack (push, ...)' at end of file">, InGroup<PragmaPack>;
-def note_pragma_pack_here : Note<
-  "previous '#pragma pack' directive that modifies alignment is here">;
 // Follow the Microsoft implementation.
 def warn_pragma_pack_show : Warning<"value of #pragma pack(show) == %0">;
 def warn_pragma_pack_pop_identifer_and_alignment : Warning<

clang/include/clang/Lex/PPCallbacks.h

@@ -381,12 +381,6 @@ public:
     Second->Ident(Loc, str);
   }
 
-  void PragmaDirective(SourceLocation Loc,
-                       PragmaIntroducerKind Introducer) override {
-    First->PragmaDirective(Loc, Introducer);
-    Second->PragmaDirective(Loc, Introducer);
-  }
-
   void PragmaComment(SourceLocation Loc, const IdentifierInfo *Kind,
                      StringRef Str) override {
     First->PragmaComment(Loc, Kind, Str);

clang/include/clang/Sema/Sema.h

@@ -208,7 +208,6 @@ namespace sema {
   class FunctionScopeInfo;
   class LambdaScopeInfo;
   class PossiblyUnreachableDiag;
-  class SemaPPCallbacks;
   class TemplateDeductionInfo;
 }
 
@@ -382,12 +381,11 @@ public:
       llvm::StringRef StackSlotLabel;
       ValueType Value;
       SourceLocation PragmaLocation;
-      SourceLocation PragmaPushLocation;
-      Slot(llvm::StringRef StackSlotLabel, ValueType Value,
-           SourceLocation PragmaLocation, SourceLocation PragmaPushLocation)
-          : StackSlotLabel(StackSlotLabel), Value(Value),
-            PragmaLocation(PragmaLocation),
-            PragmaPushLocation(PragmaPushLocation) {}
+      Slot(llvm::StringRef StackSlotLabel,
+           ValueType Value,
+           SourceLocation PragmaLocation)
+        : StackSlotLabel(StackSlotLabel), Value(Value),
+          PragmaLocation(PragmaLocation) {}
     };
     void Act(SourceLocation PragmaLocation,
              PragmaMsStackAction Action,
@@ -418,8 +416,6 @@ public:
     explicit PragmaStack(const ValueType &Default)
         : DefaultValue(Default), CurrentValue(Default) {}
 
-    bool hasValue() const { return CurrentValue != DefaultValue; }
-
     SmallVector<Slot, 2> Stack;
     ValueType DefaultValue; // Value used for PSK_Reset action.
     ValueType CurrentValue;
@@ -441,8 +437,6 @@ public:
   // Sentinel to represent when the stack is set to mac68k alignment.
   static const unsigned kMac68kAlignmentSentinel = ~0U;
   PragmaStack<unsigned> PackStack;
-  // The current #pragma pack values and locations at each #include.
-  SmallVector<std::pair<unsigned, SourceLocation>, 8> PackIncludeStack;
   // Segment #pragmas.
   PragmaStack<StringLiteral *> DataSegStack;
   PragmaStack<StringLiteral *> BSSSegStack;
@@ -8188,15 +8182,6 @@ public:
   void ActOnPragmaPack(SourceLocation PragmaLoc, PragmaMsStackAction Action,
                        StringRef SlotLabel, Expr *Alignment);
 
-  enum class PragmaPackDiagnoseKind {
-    NonDefaultStateAtInclude,
-    ChangedStateAtExit
-  };
-
-  void DiagnoseNonDefaultPragmaPack(PragmaPackDiagnoseKind Kind,
-                                    SourceLocation IncludeLoc);
-  void DiagnoseUnterminatedPragmaPack();
-
   /// ActOnPragmaMSStruct - Called on well formed \#pragma ms_struct [on|off].
   void ActOnPragmaMSStruct(PragmaMSStructKind Kind);
 
@@ -10398,12 +10383,6 @@ private:
 
   IdentifierInfo *Ident_NSError = nullptr;
 
-  /// \brief The handler for the FileChanged preprocessor events.
-  ///
-  /// Used for diagnostics that implement custom semantic analysis for #include
-  /// directives, like -Wpragma-pack.
-  sema::SemaPPCallbacks *SemaPPCallbackHandler;
-
 protected:
   friend class Parser;
   friend class InitializationSequence;

clang/include/clang/Serialization/ASTReader.h

@@ -825,7 +825,6 @@ private:
   struct PragmaPackStackEntry {
     unsigned Value;
     SourceLocation Location;
-    SourceLocation PushLocation;
     StringRef SlotLabel;
   };
   llvm::SmallVector<PragmaPackStackEntry, 2> PragmaPackStack;

clang/lib/AST/StmtProfile.cpp

@@ -1388,6 +1388,15 @@ static Stmt::StmtClass DecodeOperatorCall(const CXXOperatorCallExpr *S,
   llvm_unreachable("Invalid overloaded operator expression");
 }
 
+#if defined(_MSC_VER)
+#if _MSC_VER == 1911
+// Work around https://developercommunity.visualstudio.com/content/problem/84002/clang-cl-when-built-with-vc-2017-crashes-cause-vc.html
+// MSVC 2017 update 3 miscompiles this function, and a clang built with it
+// will crash in stage 2 of a bootstrap build.
+#pragma optimize("", off)
+#endif
+#endif
+
 void StmtProfiler::VisitCXXOperatorCallExpr(const CXXOperatorCallExpr *S) {
   if (S->isTypeDependent()) {
     // Type-dependent operator calls are profiled like their underlying
@@ -1420,6 +1429,12 @@ void StmtProfiler::VisitCXXOperatorCallExpr(const CXXOperatorCallExpr *S) {
   ID.AddInteger(S->getOperator());
 }
 
+#if defined(_MSC_VER)
+#if _MSC_VER == 1911
+#pragma optimize("", on)
+#endif
+#endif
+
 void StmtProfiler::VisitCXXMemberCallExpr(const CXXMemberCallExpr *S) {
   VisitCallExpr(S);
 }

clang/lib/AST/Type.cpp

@@ -2313,6 +2313,15 @@ bool Type::isAlignValT() const {
   return false;
 }
 
+bool Type::isStdByteType() const {
+  if (auto *ET = getAs<EnumType>()) {
+    auto *II = ET->getDecl()->getIdentifier();
+    if (II && II->isStr("byte") && ET->getDecl()->isInStdNamespace())
+      return true;
+  }
+  return false;
+}
+
 bool Type::isPromotableIntegerType() const {
   if (const BuiltinType *BT = getAs<BuiltinType>())
     switch (BT->getKind()) {

clang/lib/Basic/DiagnosticIDs.cpp

@@ -516,7 +516,7 @@ std::vector<std::string> DiagnosticIDs::getDiagnosticFlags() {
     std::string Diag(DiagGroupNames + I + 1, DiagGroupNames[I]);
     I += DiagGroupNames[I] + 1;
     Res.push_back("-W" + Diag);
-    Res.push_back("-Wno" + Diag);
+    Res.push_back("-Wno-" + Diag);
   }
 
   return Res;

clang/lib/CodeGen/CodeGenTBAA.cpp

@@ -139,6 +139,12 @@ CodeGenTBAA::getTBAAInfo(QualType QTy) {
     }
   }
 
+  // C++1z [basic.lval]p10: "If a program attempts to access the stored value of
+  // an object through a glvalue of other than one of the following types the
+  // behavior is undefined: [...] a char, unsigned char, or std::byte type."
+  if (Ty->isStdByteType())
+    return MetadataCache[Ty] = getChar();
+
   // Handle pointers.
   // TODO: Implement C++'s type "similarity" and consider dis-"similar"
   // pointers distinct.

clang/lib/Parse/ParsePragma.cpp

@@ -422,20 +422,15 @@ void Parser::HandlePragmaPack() {
   assert(Tok.is(tok::annot_pragma_pack));
   PragmaPackInfo *Info =
     static_cast<PragmaPackInfo *>(Tok.getAnnotationValue());
-  SourceLocation PragmaLoc = Tok.getLocation();
+  SourceLocation PragmaLoc = ConsumeAnnotationToken();
   ExprResult Alignment;
   if (Info->Alignment.is(tok::numeric_constant)) {
     Alignment = Actions.ActOnNumericConstant(Info->Alignment);
-    if (Alignment.isInvalid()) {
-      ConsumeAnnotationToken();
+    if (Alignment.isInvalid())
       return;
-    }
   }
   Actions.ActOnPragmaPack(PragmaLoc, Info->Action, Info->SlotLabel,
                           Alignment.get());
-  // Consume the token after processing the pragma to enable pragma-specific
-  // #include warnings.
-  ConsumeAnnotationToken();
 }
 
 void Parser::HandlePragmaMSStruct() {

clang/lib/Sema/Sema.cpp

@@ -70,49 +70,6 @@ void Sema::ActOnTranslationUnitScope(Scope *S) {
   PushDeclContext(S, Context.getTranslationUnitDecl());
 }
 
-namespace clang {
-namespace sema {
-
-class SemaPPCallbacks : public PPCallbacks {
-  Sema *S = nullptr;
-  llvm::SmallVector<SourceLocation, 8> IncludeStack;
-
-public:
-  void set(Sema &S) { this->S = &S; }
-
-  void reset() { S = nullptr; }
-
-  virtual void FileChanged(SourceLocation Loc, FileChangeReason Reason,
-                           SrcMgr::CharacteristicKind FileType,
-                           FileID PrevFID) override {
-    if (!S)
-      return;
-    switch (Reason) {
-    case EnterFile: {
-      SourceManager &SM = S->getSourceManager();
-      SourceLocation IncludeLoc = SM.getIncludeLoc(SM.getFileID(Loc));
-      if (IncludeLoc.isValid()) {
-        IncludeStack.push_back(IncludeLoc);
-        S->DiagnoseNonDefaultPragmaPack(
-            Sema::PragmaPackDiagnoseKind::NonDefaultStateAtInclude, IncludeLoc);
-      }
-      break;
-    }
-    case ExitFile:
-      if (!IncludeStack.empty())
-        S->DiagnoseNonDefaultPragmaPack(
-            Sema::PragmaPackDiagnoseKind::ChangedStateAtExit,
-            IncludeStack.pop_back_val());
-      break;
-    default:
-      break;
-    }
-  }
-};
-
-} // end namespace sema
-} // end namespace clang
-
 Sema::Sema(Preprocessor &pp, ASTContext &ctxt, ASTConsumer &consumer,
            TranslationUnitKind TUKind, CodeCompleteConsumer *CodeCompleter)
     : ExternalSource(nullptr), isMultiplexExternalSource(false),
@@ -165,12 +122,6 @@ Sema::Sema(Preprocessor &pp, ASTContext &ctxt, ASTConsumer &consumer,
 
   // Initilization of data sharing attributes stack for OpenMP
   InitDataSharingAttributesStack();
-
-  std::unique_ptr<sema::SemaPPCallbacks> Callbacks =
-      llvm::make_unique<sema::SemaPPCallbacks>();
-  SemaPPCallbackHandler = Callbacks.get();
-  PP.addPPCallbacks(std::move(Callbacks));
-  SemaPPCallbackHandler->set(*this);
 }
 
 void Sema::addImplicitTypedef(StringRef Name, QualType T) {
@@ -355,10 +306,6 @@ Sema::~Sema() {
   // Destroys data sharing attributes stack for OpenMP
   DestroyDataSharingAttributesStack();
 
-  // Detach from the PP callback handler which outlives Sema since it's owned
-  // by the preprocessor.
-  SemaPPCallbackHandler->reset();
-
   assert(DelayedTypos.empty() && "Uncorrected typos!");
 }
 
@@ -819,7 +766,6 @@ void Sema::ActOnEndOfTranslationUnit() {
     CheckDelayedMemberExceptionSpecs();
   }
 
-  DiagnoseUnterminatedPragmaPack();
   DiagnoseUnterminatedPragmaAttribute();
 
   // All delayed member exception specs should be checked or we end up accepting

clang/lib/Sema/SemaAttr.cpp

@@ -202,40 +202,6 @@ void Sema::ActOnPragmaPack(SourceLocation PragmaLoc, PragmaMsStackAction Action,
   PackStack.Act(PragmaLoc, Action, SlotLabel, AlignmentVal);
 }
 
-void Sema::DiagnoseNonDefaultPragmaPack(PragmaPackDiagnoseKind Kind,
-                                        SourceLocation IncludeLoc) {
-  if (Kind == PragmaPackDiagnoseKind::NonDefaultStateAtInclude) {
-    SourceLocation PrevLocation = PackStack.CurrentPragmaLocation;
-    // Warn about non-default alignment at #includes (without redundant
-    // warnings for the same directive in nested includes).
-    if (PackStack.hasValue() &&
-        (PackIncludeStack.empty() ||
-         PackIncludeStack.back().second != PrevLocation)) {
-      Diag(IncludeLoc, diag::warn_pragma_pack_non_default_at_include);
-      Diag(PrevLocation, diag::note_pragma_pack_here);
-    }
-    PackIncludeStack.push_back(
-        {PackStack.CurrentValue,
-         PackStack.hasValue() ? PrevLocation : SourceLocation()});
-    return;
-  }
-
-  assert(Kind == PragmaPackDiagnoseKind::ChangedStateAtExit && "invalid kind");
-  unsigned PreviousValue = PackIncludeStack.pop_back_val().first;
-  // Warn about modified alignment after #includes.
-  if (PreviousValue != PackStack.CurrentValue) {
-    Diag(IncludeLoc, diag::warn_pragma_pack_modified_after_include);
-    Diag(PackStack.CurrentPragmaLocation, diag::note_pragma_pack_here);
-  }
-}
-
-void Sema::DiagnoseUnterminatedPragmaPack() {
-  if (PackStack.Stack.empty())
-    return;
-  for (const auto &StackSlot : llvm::reverse(PackStack.Stack))
-    Diag(StackSlot.PragmaPushLocation, diag::warn_pragma_pack_no_pop_eof);
-}
-
 void Sema::ActOnPragmaMSStruct(PragmaMSStructKind Kind) { 
   MSStructPragmaOn = (Kind == PMSST_ON);
 }
@@ -283,8 +249,7 @@ void Sema::PragmaStack<ValueType>::Act(SourceLocation PragmaLocation,
     return;
   }
   if (Action & PSK_Push)
-    Stack.emplace_back(StackSlotLabel, CurrentValue, CurrentPragmaLocation,
-                       PragmaLocation);
+    Stack.push_back(Slot(StackSlotLabel, CurrentValue, CurrentPragmaLocation));
   else if (Action & PSK_Pop) {
     if (!StackSlotLabel.empty()) {
       // If we've got a label, try to find it and jump there.

clang/lib/Serialization/ASTReader.cpp

@@ -3384,7 +3384,6 @@ ASTReader::ReadASTBlock(ModuleFile &F, unsigned ClientLoadCapabilities) {
         PragmaPackStackEntry Entry;
         Entry.Value = Record[Idx++];
         Entry.Location = ReadSourceLocation(F, Record[Idx++]);
-        Entry.PushLocation = ReadSourceLocation(F, Record[Idx++]);
         PragmaPackStrings.push_back(ReadString(Record, Idx));
         Entry.SlotLabel = PragmaPackStrings.back();
         PragmaPackStack.push_back(Entry);
@@ -7580,14 +7579,13 @@ void ASTReader::UpdateSema() {
              "Expected a default alignment value");
       SemaObj->PackStack.Stack.emplace_back(
           PragmaPackStack.front().SlotLabel, SemaObj->PackStack.CurrentValue,
-          SemaObj->PackStack.CurrentPragmaLocation,
-          PragmaPackStack.front().PushLocation);
+          SemaObj->PackStack.CurrentPragmaLocation);
       DropFirst = true;
     }
     for (const auto &Entry :
          llvm::makeArrayRef(PragmaPackStack).drop_front(DropFirst ? 1 : 0))
       SemaObj->PackStack.Stack.emplace_back(Entry.SlotLabel, Entry.Value,
-                                            Entry.Location, Entry.PushLocation);
+                                            Entry.Location);
     if (PragmaPackCurrentLocation.isInvalid()) {
       assert(*PragmaPackCurrentValue == SemaObj->PackStack.DefaultValue &&
              "Expected a default alignment value");

clang/lib/Serialization/ASTWriter.cpp

@@ -4295,7 +4295,6 @@ void ASTWriter::WritePackPragmaOptions(Sema &SemaRef) {
   for (const auto &StackEntry : SemaRef.PackStack.Stack) {
     Record.push_back(StackEntry.Value);
     AddSourceLocation(StackEntry.PragmaLocation, Record);
-    AddSourceLocation(StackEntry.PragmaPushLocation, Record);
     AddString(StackEntry.StackSlotLabel, Record);
   }
   Stream.EmitRecord(PACK_PRAGMA_OPTIONS, Record);

clang/test/CodeGenCXX/std-byte.cpp

@@ -0,0 +1,41 @@
+// RUN: %clang_cc1 -std=c++1z -Werror -triple i386-unknown-unknown -emit-llvm -O1 -disable-llvm-passes -o - %s | FileCheck %s
+
+// std::byte should be considered equivalent to char for aliasing.
+
+namespace std {
+enum byte : unsigned char {};
+}
+
+// CHECK-LABEL: define void @test0(
+extern "C" void test0(std::byte *sb, int *i) {
+  // CHECK: store i8 0, i8* %{{.*}} !tbaa [[TAG_CHAR:!.*]]
+  *sb = std::byte{0};
+
+  // CHECK: store i32 1, i32* %{{.*}} !tbaa [[TAG_INT:!.*]]
+  *i = 1;
+}
+
+enum byte : unsigned char {};
+namespace my {
+enum byte : unsigned char {};
+namespace std {
+enum byte : unsigned char {};
+} // namespace std
+} // namespace my
+
+// Make sure we don't get confused with other enums named 'byte'.
+
+// CHECK-LABEL: define void @test1(
+extern "C" void test1(::byte *b, ::my::byte *mb, ::my::std::byte *msb) {
+  *b = ::byte{0};
+  *mb = ::my::byte{0};
+  *msb = ::my::std::byte{0};
+  // CHECK-NOT: store i8 0, i8* %{{.*}} !tbaa [[TAG_CHAR]]
+}
+
+// CHECK:  !"any pointer", [[TYPE_CHAR:!.*]],
+// CHECK: [[TYPE_CHAR]] = !{!"omnipotent char", [[TAG_CXX_TBAA:!.*]],
+// CHECK: [[TAG_CXX_TBAA]] = !{!"Simple C++ TBAA"}
+// CHECK: [[TAG_CHAR]] = !{[[TYPE_CHAR:!.*]], [[TYPE_CHAR]], i64 0}
+// CHECK: [[TAG_INT]] = !{[[TYPE_INT:!.*]], [[TYPE_INT]], i64 0}
+// CHECK: [[TYPE_INT]] = !{!"int", [[TYPE_CHAR]]

clang/test/Driver/autocomplete.c

@@ -42,5 +42,5 @@
 // MRELOCMODEL_CC1: -mrelocation-model
 // RUN: %clang --autocomplete=-Wma | FileCheck %s -check-prefix=WARNING
 // WARNING: -Wmacro-redefined -Wmain -Wmain-return-type -Wmalformed-warning-check -Wmany-braces-around-scalar-init -Wmax-unsigned-zero
-// RUN: %clang --autocomplete=-Wnoinvalid-pp- | FileCheck %s -check-prefix=NOWARNING
-// NOWARNING: -Wnoinvalid-pp-token
+// RUN: %clang --autocomplete=-Wno-invalid-pp- | FileCheck %s -check-prefix=NOWARNING
+// NOWARNING: -Wno-invalid-pp-token

clang/test/OpenMP/declare_simd_messages.cpp

@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -triple=x86_64-pc-win32 -verify -fopenmp -x c++ -std=c++11 -fms-extensions -Wno-pragma-pack %s
+// RUN: %clang_cc1 -triple=x86_64-pc-win32 -verify -fopenmp -x c++ -std=c++11 -fms-extensions %s
 
 // expected-error@+1 {{expected an OpenMP directive}}
 #pragma omp declare

clang/test/PCH/pragma-pack.c

@@ -1,21 +1,21 @@
 // Test this without pch.
-// RUN: %clang_cc1 -triple x86_64-apple-darwin10 %s -include %s -verify -fsyntax-only -Wno-pragma-pack -DSET
-// RUN: %clang_cc1 -triple x86_64-apple-darwin10 %s -include %s -verify -fsyntax-only -Wno-pragma-pack -DRESET
-// RUN: %clang_cc1 -triple x86_64-apple-darwin10 %s -include %s -verify -fsyntax-only -Wno-pragma-pack -DPUSH
-// RUN: %clang_cc1 -triple x86_64-apple-darwin10 %s -include %s -verify -fsyntax-only -Wno-pragma-pack -DPUSH_POP
-// RUN: %clang_cc1 -triple x86_64-apple-darwin10 %s -include %s -verify -fsyntax-only -Wno-pragma-pack -DPUSH_POP_LABEL
+// RUN: %clang_cc1 -triple x86_64-apple-darwin10 %s -include %s -verify -fsyntax-only -DSET
+// RUN: %clang_cc1 -triple x86_64-apple-darwin10 %s -include %s -verify -fsyntax-only -DRESET
+// RUN: %clang_cc1 -triple x86_64-apple-darwin10 %s -include %s -verify -fsyntax-only -DPUSH
+// RUN: %clang_cc1 -triple x86_64-apple-darwin10 %s -include %s -verify -fsyntax-only -DPUSH_POP
+// RUN: %clang_cc1 -triple x86_64-apple-darwin10 %s -include %s -verify -fsyntax-only -DPUSH_POP_LABEL
 
 // Test with pch.
-// RUN: %clang_cc1 -triple x86_64-apple-darwin10 %s -Wno-pragma-pack -DSET -emit-pch -o %t
-// RUN: %clang_cc1 -triple x86_64-apple-darwin10 %s -Wno-pragma-pack -DSET -verify -include-pch %t
-// RUN: %clang_cc1 -triple x86_64-apple-darwin10 %s -Wno-pragma-pack -DRESET -emit-pch -o %t
-// RUN: %clang_cc1 -triple x86_64-apple-darwin10 %s -Wno-pragma-pack -DRESET -verify -include-pch %t
-// RUN: %clang_cc1 -triple x86_64-apple-darwin10 %s -Wno-pragma-pack -DPUSH -emit-pch -o %t
-// RUN: %clang_cc1 -triple x86_64-apple-darwin10 %s -Wno-pragma-pack -DPUSH -verify -include-pch %t
-// RUN: %clang_cc1 -triple x86_64-apple-darwin10 %s -Wno-pragma-pack -DPUSH_POP -emit-pch -o %t
-// RUN: %clang_cc1 -triple x86_64-apple-darwin10 %s -Wno-pragma-pack -DPUSH_POP -verify -include-pch %t
-// RUN: %clang_cc1 -triple x86_64-apple-darwin10 %s -Wno-pragma-pack -DPUSH_POP_LABEL -emit-pch -o %t
-// RUN: %clang_cc1 -triple x86_64-apple-darwin10 %s -Wno-pragma-pack -DPUSH_POP_LABEL -verify -include-pch %t
+// RUN: %clang_cc1 -triple x86_64-apple-darwin10 %s -DSET -emit-pch -o %t
+// RUN: %clang_cc1 -triple x86_64-apple-darwin10 %s -DSET -verify -include-pch %t
+// RUN: %clang_cc1 -triple x86_64-apple-darwin10 %s -DRESET -emit-pch -o %t
+// RUN: %clang_cc1 -triple x86_64-apple-darwin10 %s -DRESET -verify -include-pch %t
+// RUN: %clang_cc1 -triple x86_64-apple-darwin10 %s -DPUSH -emit-pch -o %t
+// RUN: %clang_cc1 -triple x86_64-apple-darwin10 %s -DPUSH -verify -include-pch %t
+// RUN: %clang_cc1 -triple x86_64-apple-darwin10 %s -DPUSH_POP -emit-pch -o %t
+// RUN: %clang_cc1 -triple x86_64-apple-darwin10 %s -DPUSH_POP -verify -include-pch %t
+// RUN: %clang_cc1 -triple x86_64-apple-darwin10 %s -DPUSH_POP_LABEL -emit-pch -o %t
+// RUN: %clang_cc1 -triple x86_64-apple-darwin10 %s -DPUSH_POP_LABEL -verify -include-pch %t
 
 #ifndef HEADER
 #define HEADER

clang/test/PCH/suspicious-pragma-pack.c

@@ -1,8 +0,0 @@
-// RUN: %clang_cc1 -triple x86_64-apple-darwin10 %s -verify -emit-pch -o %t
-// RUN: %clang_cc1 -triple x86_64-apple-darwin10 %s -verify -include-pch %t
-
-#ifndef HEADER
-#define HEADER
-#pragma pack (push, 1)
-#endif
-// expected-warning@-2 {{unterminated '#pragma pack (push, ...)' at end of file}}

clang/test/Parser/pragma-options.c

@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -triple i386-apple-darwin9 -Wno-pragma-pack -fsyntax-only -verify %s
+// RUN: %clang_cc1 -triple i386-apple-darwin9 -fsyntax-only -verify %s
 
 /* expected-warning {{expected 'align' following '#pragma options'}} */ #pragma options
 /* expected-warning {{expected '=' following '#pragma options align'}} */ #pragma options align

clang/test/Parser/pragma-options.cpp

@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -triple i386-apple-darwin9 -Wno-pragma-pack -fsyntax-only -verify %s
+// RUN: %clang_cc1 -triple i386-apple-darwin9 -fsyntax-only -verify %s
 // expected-no-diagnostics
 
 class C {

clang/test/Parser/pragma-pack.c

@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsyntax-only -Wno-pragma-pack -verify %s
+// RUN: %clang_cc1 -fsyntax-only -verify %s
 
 // Note that this puts the expected lines before the directives to work around
 // limitations in the -verify mode.

clang/test/Sema/Inputs/pragma-pack1.h

@@ -1,23 +0,0 @@
-
-#ifdef SET_FIRST_HEADER
-#pragma pack (16)
-#ifndef SET_SECOND_HEADER
-// expected-note@-2 2 {{previous '#pragma pack' directive that modifies alignment is here}}
-#else
-// expected-note@-4 1 {{previous '#pragma pack' directive that modifies alignment is here}}
-#endif
-// expected-warning@+3 {{non-default #pragma pack value might change the alignment of struct or union members in the included file}}
-#endif
-
-#include "pragma-pack2.h"
-
-#ifdef SET_SECOND_HEADER
-// expected-warning@-3 {{the current #pragma pack aligment value is modified in the included file}}
-#endif
-
-#ifdef PUSH_POP_FIRST_HEADER
-// This is fine, we don't change the current value.
-#pragma pack (push, 4)
-
-#pragma pack (pop)
-#endif

clang/test/Sema/Inputs/pragma-pack2.h

@@ -1,6 +0,0 @@
-
-#ifdef SET_SECOND_HEADER
-#pragma pack (8) // expected-note 2 {{previous '#pragma pack' directive that modifies alignment is here}}
-#endif
-
-struct S { int x; };

clang/test/Sema/pragma-pack.c

@@ -25,8 +25,3 @@
 #pragma pack(pop, 16)
 /* expected-warning {{value of #pragma pack(show) == 16}} */ #pragma pack(show)
 
-
-// Warn about unbalanced pushes.
-#pragma pack (push,4) // expected-warning {{unterminated '#pragma pack (push, ...)' at end of file}}
-#pragma pack (push)   // expected-warning {{unterminated '#pragma pack (push, ...)' at end of file}}
-#pragma pack ()

clang/test/Sema/suspicious-pragma-pack.c

@@ -1,44 +0,0 @@
-// RUN: %clang_cc1 -triple i686-apple-darwin9 -fsyntax-only -I %S/Inputs -DSAFE -verify %s
-// RUN: %clang_cc1 -triple i686-apple-darwin9 -fsyntax-only -I %S/Inputs -DPUSH_HERE -DSAFE -verify %s
-// RUN: %clang_cc1 -triple i686-apple-darwin9 -fsyntax-only -I %S/Inputs -DPUSH_HERE -DPUSH_SET_HERE -verify %s
-// RUN: %clang_cc1 -triple i686-apple-darwin9 -fsyntax-only -I %S/Inputs -DPUSH_HERE -DRESET_HERE -DSAFE -verify %s
-// RUN: %clang_cc1 -triple i686-apple-darwin9 -fsyntax-only -I %S/Inputs -DPUSH_HERE -DSET_FIRST_HEADER -DWARN_MODIFIED_HEADER -verify %s
-// RUN: %clang_cc1 -triple i686-apple-darwin9 -fsyntax-only -I %S/Inputs -DPUSH_HERE -DRESET_HERE -DSET_FIRST_HEADER -DWARN_MODIFIED_HEADER -verify %s
-// RUN: %clang_cc1 -triple i686-apple-darwin9 -fsyntax-only -I %S/Inputs -DPUSH_HERE -DPUSH_SET_HERE -DSET_FIRST_HEADER -DWARN_MODIFIED_HEADER -verify %s
-// RUN: %clang_cc1 -triple i686-apple-darwin9 -fsyntax-only -I %S/Inputs -DPUSH_HERE -DPUSH_SET_HERE -DSET_SECOND_HEADER -DWARN_MODIFIED_HEADER -verify %s
-// RUN: %clang_cc1 -triple i686-apple-darwin9 -fsyntax-only -I %S/Inputs -DPUSH_HERE -DPUSH_SET_HERE -DSET_FIRST_HEADER -DSET_SECOND_HEADER -DWARN_MODIFIED_HEADER -verify %s
-
-// RUN: %clang_cc1 -triple i686-apple-darwin9 -fsyntax-only -I %S/Inputs -DPUSH_POP_FIRST_HEADER -DSAFE -verify %s
-
-
-#ifdef SAFE
-// expected-no-diagnostics
-#endif
-
-#ifdef PUSH_HERE
-#pragma pack (push)
-#endif
-
-#ifdef PUSH_SET_HERE
-#pragma pack (push, 4) // expected-note {{previous '#pragma pack' directive that modifies alignment is here}}
-// expected-warning@+8 {{non-default #pragma pack value might change the alignment of struct or union members in the included file}}
-#endif
-
-#ifdef RESET_HERE
-#pragma pack (4)
-#pragma pack () // no warning after reset as the value is default.
-#endif
-
-#include "pragma-pack1.h"
-
-#ifdef WARN_MODIFIED_HEADER
-// expected-warning@-3 {{the current #pragma pack aligment value is modified in the included file}}
-#endif
-
-#ifdef PUSH_SET_HERE
-#pragma pack (pop)
-#endif
-
-#ifdef PUSH_HERE
-#pragma pack (pop)
-#endif

clang/test/SemaObjC/Inputs/empty.h

@@ -1 +0,0 @@
-// empty

clang/test/SemaObjC/suspicious-pragma-pack.m

@@ -1,6 +0,0 @@
-// RUN: %clang_cc1 -triple i686-apple-darwin9 -fsyntax-only -I%S/Inputs -verify %s
-
-#pragma pack (push, 1) // expected-note {{previous '#pragma pack' directive that modifies alignment is here}}
-#import "empty.h" // expected-warning {{non-default #pragma pack value might change the alignment of struct or union members in the included file}}
-
-#pragma pack (pop)

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.c

@@ -1,40 +0,0 @@
-// RUN: %clang %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.
-void b();
-struct S {
-  int a[8];
-};
-
-int f(struct S s, unsigned i) {
-  // DEBUGGER: break 16
-  // 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 26
-    // 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 39
-  // DEBUGGER: c
-  // DEBUGGER: p x
-  // CHECK: 42
-  __block int x = 42;
-  c();
-}

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 -arch x86_64 -mllvm -fast-isel=false  %s -c -o %t.o -g
-// RUN: %clang %target_itanium_abi_host_triple -arch x86_64 %t.o -o %t.out
-// RUN: %test_debuginfo %s %t.out
-//
-// DEBUGGER: break 26
-// DEBUGGER: r
-// DEBUGGER: print mutex
-// CHECK:  ={{.*}} 0x0
-//
-// 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/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/safestack.c

@@ -1,35 +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 = ([0] = 0, [1] = 1, [2] = 2, [3] = 3, [4] = 4, [5] = 5, [6] = 6, [7] = 7)
-// DEBUGGER: c
-// DEBUGGER: p s
-// CHECK: a = ([0] = 0, [1] = 1, [2] = 2, [3] = 3, [4] = 4, [5] = 5, [6] = 6, [7] = 7)
-// DEBUGGER: c
-// DEBUGGER: p s
-// CHECK: a = ([0] = 0, [1] = 1, [2] = 2, [3] = 3, [4] = 4, [5] = 5, [6] = 6, [7] = 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)) 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/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/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/SOURCES

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

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.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/SOURCES

@@ -1,9 +0,0 @@
-math/ldexp.cl
-synchronization/barrier_impl.ll
-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/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/synchronization/barrier_impl.ll

@@ -1,32 +0,0 @@
-declare i32 @__clc_clk_local_mem_fence() #1
-declare i32 @__clc_clk_global_mem_fence() #1
-declare void @llvm.amdgcn.s.barrier() #0
-
-define void @barrier(i32 %flags) #2 {
-barrier_local_test:
-  %CLK_LOCAL_MEM_FENCE = call i32 @__clc_clk_local_mem_fence()
-  %0 = and i32 %flags, %CLK_LOCAL_MEM_FENCE
-  %1 = icmp ne i32 %0, 0
-  br i1 %1, label %barrier_local, label %barrier_global_test
-
-barrier_local:
-  call void @llvm.amdgcn.s.barrier()
-  br label %barrier_global_test
-
-barrier_global_test:
-  %CLK_GLOBAL_MEM_FENCE = call i32 @__clc_clk_global_mem_fence()
-  %2 = and i32 %flags, %CLK_GLOBAL_MEM_FENCE
-  %3 = icmp ne i32 %2, 0
-  br i1 %3, label %barrier_global, label %done
-
-barrier_global:
-  call void @llvm.amdgcn.s.barrier()
-  br label %done
-
-done:
-  ret void
-}
-
-attributes #0 = { nounwind convergent }
-attributes #1 = { nounwind alwaysinline }
-attributes #2 = { nounwind convergent alwaysinline }

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.ll

@@ -1,21 +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
-
-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.ll

@@ -1,21 +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
-
-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.ll

@@ -1,21 +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
-
-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()
-{
-	__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,18 +0,0 @@
-atomic/atomic.cl
-math/nextafter.cl
-math/sqrt.cl
-synchronization/barrier.cl
-image/get_image_width.cl
-image/get_image_height.cl
-image/get_image_depth.cl
-image/get_image_channel_data_type.cl
-image/get_image_channel_order.cl
-image/get_image_attributes_impl.ll
-image/read_imagef.cl
-image/read_imagei.cl
-image/read_imageui.cl
-image/read_image_impl.ll
-image/write_imagef.cl
-image/write_imagei.cl
-image/write_imageui.cl
-image/write_image_impl.ll

libclc/amdgpu/lib/atomic/atomic.cl

@@ -1,65 +0,0 @@
-#include <clc/clc.h>
-
-#define ATOMIC_FUNC_DEFINE(RET_SIGN, ARG_SIGN, TYPE, CL_FUNCTION, CLC_FUNCTION, CL_ADDRSPACE, LLVM_ADDRSPACE) \
-_CLC_OVERLOAD _CLC_DEF RET_SIGN TYPE CL_FUNCTION (volatile CL_ADDRSPACE RET_SIGN TYPE *p, RET_SIGN TYPE val) { \
-	return (RET_SIGN TYPE)__clc_##CLC_FUNCTION##_addr##LLVM_ADDRSPACE((volatile CL_ADDRSPACE ARG_SIGN TYPE*)p, (ARG_SIGN TYPE)val); \
-}
-
-/* For atomic functions that don't need different bitcode dependending on argument signedness */
-#define ATOMIC_FUNC_SIGN(TYPE, FUNCTION, CL_ADDRSPACE, LLVM_ADDRSPACE) \
-	_CLC_DECL signed TYPE __clc_##FUNCTION##_addr##LLVM_ADDRSPACE(volatile CL_ADDRSPACE signed TYPE*, signed TYPE); \
-	ATOMIC_FUNC_DEFINE(signed, signed, TYPE, FUNCTION, FUNCTION, CL_ADDRSPACE, LLVM_ADDRSPACE) \
-	ATOMIC_FUNC_DEFINE(unsigned, signed, TYPE, FUNCTION, FUNCTION, CL_ADDRSPACE, LLVM_ADDRSPACE)
-
-#define ATOMIC_FUNC_ADDRSPACE(TYPE, FUNCTION) \
-	ATOMIC_FUNC_SIGN(TYPE, FUNCTION, global, 1) \
-	ATOMIC_FUNC_SIGN(TYPE, FUNCTION, local, 3)
-
-#define ATOMIC_FUNC(FUNCTION) \
-	ATOMIC_FUNC_ADDRSPACE(int, FUNCTION)
-
-#define ATOMIC_FUNC_DEFINE_3_ARG(RET_SIGN, ARG_SIGN, TYPE, CL_FUNCTION, CLC_FUNCTION, CL_ADDRSPACE, LLVM_ADDRSPACE) \
-_CLC_OVERLOAD _CLC_DEF RET_SIGN TYPE CL_FUNCTION (volatile CL_ADDRSPACE RET_SIGN TYPE *p, RET_SIGN TYPE cmp, RET_SIGN TYPE val) { \
-	return (RET_SIGN TYPE)__clc_##CLC_FUNCTION##_addr##LLVM_ADDRSPACE((volatile CL_ADDRSPACE ARG_SIGN TYPE*)p, (ARG_SIGN TYPE)cmp, (ARG_SIGN TYPE)val); \
-}
-
-/* For atomic functions that don't need different bitcode dependending on argument signedness */
-#define ATOMIC_FUNC_SIGN_3_ARG(TYPE, FUNCTION, CL_ADDRSPACE, LLVM_ADDRSPACE) \
-	_CLC_DECL signed TYPE __clc_##FUNCTION##_addr##LLVM_ADDRSPACE(volatile CL_ADDRSPACE signed TYPE*, signed TYPE, signed TYPE); \
-	ATOMIC_FUNC_DEFINE_3_ARG(signed, signed, TYPE, FUNCTION, FUNCTION, CL_ADDRSPACE, LLVM_ADDRSPACE) \
-	ATOMIC_FUNC_DEFINE_3_ARG(unsigned, signed, TYPE, FUNCTION, FUNCTION, CL_ADDRSPACE, LLVM_ADDRSPACE)
-
-#define ATOMIC_FUNC_ADDRSPACE_3_ARG(TYPE, FUNCTION) \
-	ATOMIC_FUNC_SIGN_3_ARG(TYPE, FUNCTION, global, 1) \
-	ATOMIC_FUNC_SIGN_3_ARG(TYPE, FUNCTION, local, 3)
-
-#define ATOMIC_FUNC_3_ARG(FUNCTION) \
-	ATOMIC_FUNC_ADDRSPACE_3_ARG(int, FUNCTION)
-
-ATOMIC_FUNC(atomic_add)
-ATOMIC_FUNC(atomic_and)
-ATOMIC_FUNC(atomic_or)
-ATOMIC_FUNC(atomic_sub)
-ATOMIC_FUNC(atomic_xchg)
-ATOMIC_FUNC(atomic_xor)
-ATOMIC_FUNC_3_ARG(atomic_cmpxchg)
-
-_CLC_DECL signed int __clc_atomic_max_addr1(volatile global signed int*, signed int);
-_CLC_DECL signed int __clc_atomic_max_addr3(volatile local signed int*, signed int);
-_CLC_DECL uint __clc_atomic_umax_addr1(volatile global uint*, uint);
-_CLC_DECL uint __clc_atomic_umax_addr3(volatile local uint*, uint);
-
-ATOMIC_FUNC_DEFINE(signed, signed, int, atomic_max, atomic_max, global, 1)
-ATOMIC_FUNC_DEFINE(signed, signed, int, atomic_max, atomic_max, local, 3)
-ATOMIC_FUNC_DEFINE(unsigned, unsigned, int, atomic_max, atomic_umax, global, 1)
-ATOMIC_FUNC_DEFINE(unsigned, unsigned, int, atomic_max, atomic_umax, local, 3)
-
-_CLC_DECL signed int __clc_atomic_min_addr1(volatile global signed int*, signed int);
-_CLC_DECL signed int __clc_atomic_min_addr3(volatile local signed int*, signed int);
-_CLC_DECL uint __clc_atomic_umin_addr1(volatile global uint*, uint);
-_CLC_DECL uint __clc_atomic_umin_addr3(volatile local uint*, uint);
-
-ATOMIC_FUNC_DEFINE(signed, signed, int, atomic_min, atomic_min, global, 1)
-ATOMIC_FUNC_DEFINE(signed, signed, int, atomic_min, atomic_min, local, 3)
-ATOMIC_FUNC_DEFINE(unsigned, unsigned, int, atomic_min, atomic_umin, global, 1)
-ATOMIC_FUNC_DEFINE(unsigned, unsigned, int, atomic_min, atomic_umin, local, 3)

libclc/amdgpu/lib/image/get_image_attributes_impl.ll

@@ -1,87 +0,0 @@
-%opencl.image2d_t = type opaque
-%opencl.image3d_t = type opaque
-
-declare i32 @llvm.OpenCL.image.get.resource.id.2d(
-  %opencl.image2d_t addrspace(1)*) nounwind readnone
-declare i32 @llvm.OpenCL.image.get.resource.id.3d(
-  %opencl.image3d_t addrspace(1)*) nounwind readnone
-
-declare [3 x i32] @llvm.OpenCL.image.get.size.2d(
-  %opencl.image2d_t addrspace(1)*) nounwind readnone
-declare [3 x i32] @llvm.OpenCL.image.get.size.3d(
-  %opencl.image3d_t addrspace(1)*) nounwind readnone
-
-declare [2 x i32] @llvm.OpenCL.image.get.format.2d(
-  %opencl.image2d_t addrspace(1)*) nounwind readnone
-declare [2 x i32] @llvm.OpenCL.image.get.format.3d(
-  %opencl.image3d_t addrspace(1)*) nounwind readnone
-
-define i32 @__clc_get_image_width_2d(
-                          %opencl.image2d_t addrspace(1)* nocapture %img) #0 {
-  %1 = tail call [3 x i32] @llvm.OpenCL.image.get.size.2d(
-    %opencl.image2d_t addrspace(1)* %img)
-  %2 = extractvalue [3 x i32] %1, 0
-  ret i32 %2
-}
-define i32 @__clc_get_image_width_3d(
-                          %opencl.image3d_t addrspace(1)* nocapture %img) #0 {
-  %1 = tail call [3 x i32] @llvm.OpenCL.image.get.size.3d(
-    %opencl.image3d_t addrspace(1)* %img)
-  %2 = extractvalue [3 x i32] %1, 0
-  ret i32 %2
-}
-
-define i32 @__clc_get_image_height_2d(
-                          %opencl.image2d_t addrspace(1)* nocapture %img) #0 {
-  %1 = tail call [3 x i32] @llvm.OpenCL.image.get.size.2d(
-    %opencl.image2d_t addrspace(1)* %img)
-  %2 = extractvalue [3 x i32] %1, 1
-  ret i32 %2
-}
-define i32 @__clc_get_image_height_3d(
-                          %opencl.image3d_t addrspace(1)* nocapture %img) #0 {
-  %1 = tail call [3 x i32] @llvm.OpenCL.image.get.size.3d(
-    %opencl.image3d_t addrspace(1)* %img)
-  %2 = extractvalue [3 x i32] %1, 1
-  ret i32 %2
-}
-
-define i32 @__clc_get_image_depth_3d(
-                          %opencl.image3d_t addrspace(1)* nocapture %img) #0 {
-  %1 = tail call [3 x i32] @llvm.OpenCL.image.get.size.3d(
-    %opencl.image3d_t addrspace(1)* %img)
-  %2 = extractvalue [3 x i32] %1, 2
-  ret i32 %2
-}
-
-define i32 @__clc_get_image_channel_data_type_2d(
-                          %opencl.image2d_t addrspace(1)* nocapture %img) #0 {
-  %1 = tail call [2 x i32] @llvm.OpenCL.image.get.format.2d(
-    %opencl.image2d_t addrspace(1)* %img)
-  %2 = extractvalue [2 x i32] %1, 0
-  ret i32 %2
-}
-define i32 @__clc_get_image_channel_data_type_3d(
-                          %opencl.image3d_t addrspace(1)* nocapture %img) #0 {
-  %1 = tail call [2 x i32] @llvm.OpenCL.image.get.format.3d(
-    %opencl.image3d_t addrspace(1)* %img)
-  %2 = extractvalue [2 x i32] %1, 0
-  ret i32 %2
-}
-
-define i32 @__clc_get_image_channel_order_2d(
-                          %opencl.image2d_t addrspace(1)* nocapture %img) #0 {
-  %1 = tail call [2 x i32] @llvm.OpenCL.image.get.format.2d(
-    %opencl.image2d_t addrspace(1)* %img)
-  %2 = extractvalue [2 x i32] %1, 1
-  ret i32 %2
-}
-define i32 @__clc_get_image_channel_order_3d(
-                          %opencl.image3d_t addrspace(1)* nocapture %img) #0 {
-  %1 = tail call [2 x i32] @llvm.OpenCL.image.get.format.3d(
-    %opencl.image3d_t addrspace(1)* %img)
-  %2 = extractvalue [2 x i32] %1, 1
-  ret i32 %2
-}
-
-attributes #0 = { nounwind readnone alwaysinline }

libclc/amdgpu/lib/image/get_image_channel_data_type.cl

@@ -1,13 +0,0 @@
-#include <clc/clc.h>
-
-_CLC_DECL int __clc_get_image_channel_data_type_2d(image2d_t);
-_CLC_DECL int __clc_get_image_channel_data_type_3d(image3d_t);
-
-_CLC_OVERLOAD _CLC_DEF int
-get_image_channel_data_type(image2d_t image) {
-  return __clc_get_image_channel_data_type_2d(image);
-}
-_CLC_OVERLOAD _CLC_DEF int
-get_image_channel_data_type(image3d_t image) {
-  return __clc_get_image_channel_data_type_3d(image);
-}

libclc/amdgpu/lib/image/get_image_channel_order.cl

@@ -1,13 +0,0 @@
-#include <clc/clc.h>
-
-_CLC_DECL int __clc_get_image_channel_order_2d(image2d_t);
-_CLC_DECL int __clc_get_image_channel_order_3d(image3d_t);
-
-_CLC_OVERLOAD _CLC_DEF int
-get_image_channel_order(image2d_t image) {
-  return __clc_get_image_channel_order_2d(image);
-}
-_CLC_OVERLOAD _CLC_DEF int
-get_image_channel_order(image3d_t image) {
-  return __clc_get_image_channel_order_3d(image);
-}

libclc/amdgpu/lib/image/get_image_depth.cl

@@ -1,8 +0,0 @@
-#include <clc/clc.h>
-
-_CLC_DECL int __clc_get_image_depth_3d(image3d_t);
-
-_CLC_OVERLOAD _CLC_DEF int
-get_image_depth(image3d_t image) {
-	return __clc_get_image_depth_3d(image);
-}

libclc/amdgpu/lib/image/get_image_height.cl

@@ -1,13 +0,0 @@
-#include <clc/clc.h>
-
-_CLC_DECL int __clc_get_image_height_2d(image2d_t);
-_CLC_DECL int __clc_get_image_height_3d(image3d_t);
-
-_CLC_OVERLOAD _CLC_DEF int
-get_image_height(image2d_t image) {
-  return __clc_get_image_height_2d(image);
-}
-_CLC_OVERLOAD _CLC_DEF int
-get_image_height(image3d_t image) {
-  return __clc_get_image_height_3d(image);
-}

libclc/amdgpu/lib/image/get_image_width.cl

@@ -1,13 +0,0 @@
-#include <clc/clc.h>
-
-_CLC_DECL int __clc_get_image_width_2d(image2d_t);
-_CLC_DECL int __clc_get_image_width_3d(image3d_t);
-
-_CLC_OVERLOAD _CLC_DEF int
-get_image_width(image2d_t image) {
-  return __clc_get_image_width_2d(image);
-}
-_CLC_OVERLOAD _CLC_DEF int
-get_image_width(image3d_t image) {
-  return __clc_get_image_width_3d(image);
-}

libclc/amdgpu/lib/image/read_image_impl.ll

@@ -1,46 +0,0 @@
-%opencl.image2d_t = type opaque
-
-declare <4 x float> @llvm.R600.tex(<4 x float>, i32, i32, i32, i32, i32, i32,
-                                   i32, i32, i32) readnone
-declare i32 @llvm.OpenCL.image.get.resource.id.2d(
-  %opencl.image2d_t addrspace(1)*) nounwind readnone
-declare i32 @llvm.OpenCL.sampler.get.resource.id(i32) readnone
-
-define <4 x float> @__clc_v4f_from_v2f(<2 x float> %v) alwaysinline {
-  %e0 = extractelement <2 x float> %v, i32 0
-  %e1 = extractelement <2 x float> %v, i32 1
-  %res.0 = insertelement <4 x float> undef,  float %e0, i32 0
-  %res.1 = insertelement <4 x float> %res.0, float %e1, i32 1
-  %res.2 = insertelement <4 x float> %res.1, float 0.0, i32 2
-  %res.3 = insertelement <4 x float> %res.2, float 0.0, i32 3
-  ret <4 x float> %res.3
-}
-
-define <4 x float> @__clc_read_imagef_tex(
-    %opencl.image2d_t addrspace(1)* nocapture %img,
-    i32 %sampler, <2 x float> %coord) alwaysinline {
-entry:
-  %coord_v4 = call <4 x float> @__clc_v4f_from_v2f(<2 x float> %coord)
-  %smp_id = call i32 @llvm.OpenCL.sampler.get.resource.id(i32 %sampler)
-  %img_id = call i32 @llvm.OpenCL.image.get.resource.id.2d(
-      %opencl.image2d_t addrspace(1)* %img)
-  %tex_id = add i32 %img_id, 2    ; First 2 IDs are reserved.
-
-  %coord_norm = and i32 %sampler, 1
-  %is_norm = icmp eq i32 %coord_norm, 1
-  br i1 %is_norm, label %NormCoord, label %UnnormCoord
-NormCoord:
-  %data.norm = call <4 x float> @llvm.R600.tex(
-      <4 x float> %coord_v4,
-      i32 0, i32 0, i32 0,        ; Offset.
-      i32 2, i32 %smp_id,
-      i32 1, i32 1, i32 1, i32 1) ; Normalized coords.
-  ret <4 x float> %data.norm
-UnnormCoord:
-  %data.unnorm = call <4 x float> @llvm.R600.tex(
-      <4 x float> %coord_v4,
-      i32 0, i32 0, i32 0,        ; Offset.
-      i32 %tex_id, i32 %smp_id,
-      i32 0, i32 0, i32 0, i32 0) ; Unnormalized coords.
-  ret <4 x float> %data.unnorm
-}

libclc/amdgpu/lib/image/read_imagef.cl

@@ -1,14 +0,0 @@
-#include <clc/clc.h>
-
-_CLC_DECL float4 __clc_read_imagef_tex(image2d_t, sampler_t, float2);
-
-_CLC_OVERLOAD _CLC_DEF float4 read_imagef(image2d_t image, sampler_t sampler,
-                                          int2 coord) {
-  float2 coord_float = (float2)(coord.x, coord.y);
-  return __clc_read_imagef_tex(image, sampler, coord_float);
-}
-
-_CLC_OVERLOAD _CLC_DEF float4 read_imagef(image2d_t image, sampler_t sampler,
-                                          float2 coord) {
-  return __clc_read_imagef_tex(image, sampler, coord);
-}

libclc/amdgpu/lib/image/read_imagei.cl

@@ -1,23 +0,0 @@
-#include <clc/clc.h>
-
-_CLC_DECL float4 __clc_read_imagef_tex(image2d_t, sampler_t, float2);
-
-int4 __clc_reinterpret_v4f_to_v4i(float4 v) {
-  union {
-    int4 v4i;
-    float4 v4f;
-  } res = { .v4f = v};
-  return res.v4i;
-}
-
-_CLC_OVERLOAD _CLC_DEF int4 read_imagei(image2d_t image, sampler_t sampler,
-                                        int2 coord) {
-  float2 coord_float = (float2)(coord.x, coord.y);
-  return __clc_reinterpret_v4f_to_v4i(
-    __clc_read_imagef_tex(image, sampler, coord_float));
-}
-_CLC_OVERLOAD _CLC_DEF int4 read_imagei(image2d_t image, sampler_t sampler,
-                                        float2 coord) {
-  return __clc_reinterpret_v4f_to_v4i(
-    __clc_read_imagef_tex(image, sampler, coord));
-}

libclc/amdgpu/lib/image/read_imageui.cl

@@ -1,23 +0,0 @@
-#include <clc/clc.h>
-
-_CLC_DECL float4 __clc_read_imagef_tex(image2d_t, sampler_t, float2);
-
-uint4 __clc_reinterpret_v4f_to_v4ui(float4 v) {
-  union {
-    uint4 v4ui;
-    float4 v4f;
-  } res = { .v4f = v};
-  return res.v4ui;
-}
-
-_CLC_OVERLOAD _CLC_DEF uint4 read_imageui(image2d_t image, sampler_t sampler,
-                                          int2 coord) {
-  float2 coord_float = (float2)(coord.x, coord.y);
-  return __clc_reinterpret_v4f_to_v4ui(
-    __clc_read_imagef_tex(image, sampler, coord_float));
-}
-_CLC_OVERLOAD _CLC_DEF uint4 read_imageui(image2d_t image, sampler_t sampler,
-                                          float2 coord) {
-  return __clc_reinterpret_v4f_to_v4ui(
-    __clc_read_imagef_tex(image, sampler, coord));
-}

libclc/amdgpu/lib/image/write_image_impl.ll

@@ -1,52 +0,0 @@
-%opencl.image2d_t = type opaque
-%opencl.image3d_t = type opaque
-
-declare i32 @llvm.OpenCL.image.get.resource.id.2d(
-  %opencl.image2d_t addrspace(1)*) nounwind readnone
-declare i32 @llvm.OpenCL.image.get.resource.id.3d(
-  %opencl.image3d_t addrspace(1)*) nounwind readnone
-
-declare void @llvm.r600.rat.store.typed(<4 x i32> %color, <4 x i32> %coord, i32 %rat_id)
-
-define void @__clc_write_imageui_2d(
-    %opencl.image2d_t addrspace(1)* nocapture %img,
-    <2 x i32> %coord, <4 x i32> %color) #0 {
-
-  ; Coordinate int2 -> int4.
-  %e0 = extractelement <2 x i32> %coord, i32 0
-  %e1 = extractelement <2 x i32> %coord, i32 1
-  %coord.0 = insertelement <4 x i32> undef,    i32 %e0, i32 0
-  %coord.1 = insertelement <4 x i32> %coord.0, i32 %e1, i32 1
-  %coord.2 = insertelement <4 x i32> %coord.1, i32 0,  i32 2
-  %coord.3 = insertelement <4 x i32> %coord.2, i32 0,  i32 3
-
-  ; Get RAT ID.
-  %img_id = call i32 @llvm.OpenCL.image.get.resource.id.2d(
-      %opencl.image2d_t addrspace(1)* %img)
-  %rat_id = add i32 %img_id, 1
-
-  ; Call store intrinsic.
-  call void @llvm.r600.rat.store.typed(<4 x i32> %color, <4 x i32> %coord.3, i32 %rat_id)
-  ret void
-}
-
-define void @__clc_write_imagei_2d(
-    %opencl.image2d_t addrspace(1)* nocapture %img,
-    <2 x i32> %coord, <4 x i32> %color) #0 {
-  call void @__clc_write_imageui_2d(
-      %opencl.image2d_t addrspace(1)* nocapture %img,
-      <2 x i32> %coord, <4 x i32> %color)
-  ret void
-}
-
-define void @__clc_write_imagef_2d(
-    %opencl.image2d_t addrspace(1)* nocapture %img,
-    <2 x i32> %coord, <4 x float> %color) #0 {
-  %color.i32 = bitcast <4 x float> %color to <4 x i32>
-  call void @__clc_write_imageui_2d(
-      %opencl.image2d_t addrspace(1)* nocapture %img,
-      <2 x i32> %coord, <4 x i32> %color.i32)
-  ret void
-}
-
-attributes #0 = { alwaysinline }

libclc/amdgpu/lib/image/write_imagef.cl

@@ -1,9 +0,0 @@
-#include <clc/clc.h>
-
-_CLC_DECL void __clc_write_imagef_2d(image2d_t image, int2 coord, float4 color);
-
-_CLC_OVERLOAD _CLC_DEF void
-write_imagef(image2d_t image, int2 coord, float4 color)
-{
-  __clc_write_imagef_2d(image, coord, color);
-}

libclc/amdgpu/lib/image/write_imagei.cl

@@ -1,9 +0,0 @@
-#include <clc/clc.h>
-
-_CLC_DECL void __clc_write_imagei_2d(image2d_t image, int2 coord, int4 color);
-
-_CLC_OVERLOAD _CLC_DEF void
-write_imagei(image2d_t image, int2 coord, int4 color)
-{
-  __clc_write_imagei_2d(image, coord, color);
-}

libclc/amdgpu/lib/image/write_imageui.cl

@@ -1,9 +0,0 @@
-#include <clc/clc.h>
-
-_CLC_DECL void __clc_write_imageui_2d(image2d_t image, int2 coord, uint4 color);
-
-_CLC_OVERLOAD _CLC_DEF void
-write_imageui(image2d_t image, int2 coord, uint4 color)
-{
-  __clc_write_imageui_2d(image, coord, color);
-}

libclc/amdgpu/lib/math/nextafter.cl

@@ -1,9 +0,0 @@
-#include <clc/clc.h>
-#include "../lib/clcmacro.h"
-
-_CLC_DEFINE_BINARY_BUILTIN(float, nextafter, __clc_nextafter, float, float)
-
-#ifdef cl_khr_fp64
-#pragma OPENCL EXTENSION cl_khr_fp64 : enable
-_CLC_DEFINE_BINARY_BUILTIN(double, nextafter, __clc_nextafter, double, double)
-#endif

libclc/amdgpu/lib/math/sqrt.cl

@@ -1,64 +0,0 @@
-/*
- * Copyright (c) 2015 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"
-#include "math/clc_sqrt.h"
-
-_CLC_DEFINE_UNARY_BUILTIN(float, sqrt, __clc_sqrt, float)
-
-#ifdef cl_khr_fp64
-
-#pragma OPENCL EXTENSION cl_khr_fp64 : enable
-
-#ifdef __AMDGCN__
-  #define __clc_builtin_rsq __builtin_amdgcn_rsq
-#else
-  #define __clc_builtin_rsq __builtin_r600_recipsqrt_ieee
-#endif
-
-_CLC_OVERLOAD _CLC_DEF double sqrt(double x) {
-
-  uint vcc = x < 0x1p-767;
-  uint exp0 = vcc ? 0x100 : 0;
-  unsigned exp1 = vcc ? 0xffffff80 : 0;
-
-  double v01 = ldexp(x, exp0);
-  double v23 = __clc_builtin_rsq(v01);
-  double v45 = v01 * v23;
-  v23 = v23 * 0.5;
-
-  double v67 = fma(-v23, v45, 0.5);
-  v45 = fma(v45, v67, v45);
-  double v89 = fma(-v45, v45, v01);
-  v23 = fma(v23, v67, v23);
-  v45 = fma(v89, v23, v45);
-  v67 = fma(-v45, v45, v01);
-  v23 = fma(v67, v23, v45);
-
-  v23 = ldexp(v23, exp1);
-  return ((x == __builtin_inf()) || (x == 0.0)) ? v01 : v23;
-}
-
-_CLC_UNARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, double, sqrt, double);
-
-#endif

libclc/amdgpu/lib/synchronization/barrier.cl

@@ -1,10 +0,0 @@
-
-#include <clc/clc.h>
-
-_CLC_DEF int __clc_clk_local_mem_fence() {
-  return CLK_LOCAL_MEM_FENCE;
-}
-
-_CLC_DEF int __clc_clk_global_mem_fence() {
-  return CLK_GLOBAL_MEM_FENCE;
-}

libclc/build/metabuild.py

@@ -1,100 +0,0 @@
-import ninja_syntax
-import os
-
-# Simple meta-build system.
-
-class Make(object):
-  def __init__(self):
-    self.output = open(self.output_filename(), 'w')
-    self.rules = {}
-    self.rule_text = ''
-    self.all_targets = []
-    self.default_targets = []
-    self.clean_files = []
-    self.distclean_files = []
-    self.output.write("""all::
-
-ifndef VERBOSE
-  Verb = @
-endif
-
-""")
-
-  def output_filename(self):
-    return 'Makefile'
-
-  def rule(self, name, command, description=None, depfile=None,
-           generator=False):
-    self.rules[name] = {'command': command, 'description': description,
-                        'depfile': depfile, 'generator': generator}
-
-  def build(self, output, rule, inputs=[], implicit=[], order_only=[]):
-    inputs = self._as_list(inputs)
-    implicit = self._as_list(implicit)
-    order_only = self._as_list(order_only)
-
-    output_dir = os.path.dirname(output)
-    if output_dir != '' and not os.path.isdir(output_dir):
-      os.makedirs(output_dir)
-
-    dollar_in = ' '.join(inputs)
-    subst = lambda text: text.replace('$in', dollar_in).replace('$out', output)
-
-    deps = ' '.join(inputs + implicit)
-    if order_only:
-      deps += ' | '
-      deps += ' '.join(order_only)
-    self.output.write('%s: %s\n' % (output, deps))
-
-    r = self.rules[rule]
-    command = subst(r['command'])
-    if r['description']:
-      desc = subst(r['description'])
-      self.output.write('\t@echo %s\n\t$(Verb) %s\n' % (desc, command))
-    else:
-      self.output.write('\t%s\n' % command)
-    if r['depfile']:
-      depfile = subst(r['depfile'])
-      self.output.write('-include '+depfile+'\n')
-    self.output.write('\n')
-
-    self.all_targets.append(output)
-    if r['generator']:
-      self.distclean_files.append(output)
-      if r['depfile']:
-        self.distclean_files.append(depfile)
-    else:
-      self.clean_files.append(output)
-      if r['depfile']:
-        self.distclean_files.append(depfile)
-
-
-  def _as_list(self, input):
-    if isinstance(input, list):
-      return input
-    return [input]
-
-  def default(self, paths):
-    self.default_targets += self._as_list(paths)
-
-  def finish(self):
-    self.output.write('all:: %s\n\n' % ' '.join(self.default_targets or self.all_targets))
-    self.output.write('clean: \n\trm -f %s\n\n' % ' '.join(self.clean_files))
-    self.output.write('distclean: clean\n\trm -f %s\n' % ' '.join(self.distclean_files))
-
-class Ninja(ninja_syntax.Writer):
-  def __init__(self):
-    ninja_syntax.Writer.__init__(self, open(self.output_filename(), 'w'))
-
-  def output_filename(self):
-    return 'build.ninja'
-
-  def finish(self):
-    pass
-
-def from_name(name):
-  if name == 'make':
-    return Make()
-  if name == 'ninja':
-    return Ninja()
-  raise LookupError, 'unknown generator: %s; supported generators are make and ninja' % name

libclc/build/ninja_syntax.py

@@ -1,118 +0,0 @@
-#!/usr/bin/python
-
-"""Python module for generating .ninja files.
-
-Note that this is emphatically not a required piece of Ninja; it's
-just a helpful utility for build-file-generation systems that already
-use Python.
-"""
-
-import textwrap
-import re
-
-class Writer(object):
-    def __init__(self, output, width=78):
-        self.output = output
-        self.width = width
-
-    def newline(self):
-        self.output.write('\n')
-
-    def comment(self, text):
-        for line in textwrap.wrap(text, self.width - 2):
-            self.output.write('# ' + line + '\n')
-
-    def variable(self, key, value, indent=0):
-        if value is None:
-            return
-        if isinstance(value, list):
-            value = ' '.join(value)
-        self._line('%s = %s' % (key, value), indent)
-
-    def rule(self, name, command, description=None, depfile=None,
-             generator=False):
-        self._line('rule %s' % name)
-        self.variable('command', escape(command), indent=1)
-        if description:
-            self.variable('description', description, indent=1)
-        if depfile:
-            self.variable('depfile', depfile, indent=1)
-        if generator:
-            self.variable('generator', '1', indent=1)
-
-    def build(self, outputs, rule, inputs=None, implicit=None, order_only=None,
-              variables=None):
-        outputs = self._as_list(outputs)
-        all_inputs = self._as_list(inputs)[:]
-
-        if implicit:
-            all_inputs.append('|')
-            all_inputs.extend(self._as_list(implicit))
-        if order_only:
-            all_inputs.append('||')
-            all_inputs.extend(self._as_list(order_only))
-
-        self._line('build %s: %s %s' % (' '.join(outputs),
-                                        rule,
-                                        ' '.join(all_inputs)))
-
-        if variables:
-            for key, val in variables:
-                self.variable(key, val, indent=1)
-
-        return outputs
-
-    def include(self, path):
-        self._line('include %s' % path)
-
-    def subninja(self, path):
-        self._line('subninja %s' % path)
-
-    def default(self, paths):
-        self._line('default %s' % ' '.join(self._as_list(paths)))
-
-    def _line(self, text, indent=0):
-        """Write 'text' word-wrapped at self.width characters."""
-        leading_space = '  ' * indent
-        while len(text) > self.width:
-            # The text is too wide; wrap if possible.
-
-            # Find the rightmost space that would obey our width constraint.
-            available_space = self.width - len(leading_space) - len(' $')
-            space = text.rfind(' ', 0, available_space)
-            if space < 0:
-                # No such space; just use the first space we can find.
-                space = text.find(' ', available_space)
-            if space < 0:
-                # Give up on breaking.
-                break
-
-            self.output.write(leading_space + text[0:space] + ' $\n')
-            text = text[space+1:]
-
-            # Subsequent lines are continuations, so indent them.
-            leading_space = '  ' * (indent+2)
-
-        self.output.write(leading_space + text + '\n')
-
-    def _as_list(self, input):
-        if input is None:
-            return []
-        if isinstance(input, list):
-            return input
-        return [input]
-
-
-def escape(string):
-    """Escape a string such that Makefile and shell variables are
-       correctly escaped for use in a Ninja file.
-    """
-    assert '\n' not in string, 'Ninja syntax does not allow newlines'
-    # We only have one special metacharacter: '$'.
-
-    # We should leave $in and $out untouched.
-    # Just look for makefile/shell style substitutions
-    return re.sub(r'(\$[{(][a-z_]+[})])',
-                  r'$\1',
-                  string,
-                  flags=re.IGNORECASE)

libclc/compile-test.sh

@@ -1,3 +0,0 @@
-#!/bin/sh
-
-clang -target nvptx--nvidiacl -Iptx-nvidiacl/include -Igeneric/include -Xclang -mlink-bitcode-file -Xclang nvptx--nvidiacl/lib/builtins.bc -include clc/clc.h -Dcl_clang_storage_class_specifiers -Dcl_khr_fp64 "$@"

libclc/configure.py

@@ -1,286 +0,0 @@
-#!/usr/bin/python
-
-def c_compiler_rule(b, name, description, compiler, flags):
-  command = "%s -MMD -MF $out.d %s -c -o $out $in" % (compiler, flags)
-  b.rule(name, command, description + " $out", depfile="$out.d")
-
-version_major = 0;
-version_minor = 2;
-version_patch = 0;
-
-from optparse import OptionParser
-import os
-import string
-from subprocess import *
-import sys
-
-srcdir = os.path.dirname(sys.argv[0])
-
-sys.path.insert(0, os.path.join(srcdir, 'build'))
-import metabuild
-
-p = OptionParser()
-p.add_option('--with-llvm-config', metavar='PATH',
-             help='use given llvm-config script')
-p.add_option('--with-cxx-compiler', metavar='PATH',
-             help='use given C++ compiler')
-p.add_option('--prefix', metavar='PATH',
-             help='install to given prefix')
-p.add_option('--libexecdir', metavar='PATH',
-             help='install *.bc to given dir')
-p.add_option('--includedir', metavar='PATH',
-             help='install include files to given dir')
-p.add_option('--pkgconfigdir', metavar='PATH',
-             help='install clc.pc to given dir')
-p.add_option('-g', metavar='GENERATOR', default='make',
-             help='use given generator (default: make)')
-p.add_option('--enable-runtime-subnormal', action="store_true", default=False,
-             help='Allow runtimes to choose subnormal support')
-(options, args) = p.parse_args()
-
-llvm_config_exe = options.with_llvm_config or "llvm-config"
-
-prefix = options.prefix
-if not prefix:
-  prefix = '/usr/local'
-
-libexecdir = options.libexecdir
-if not libexecdir:
-  libexecdir = os.path.join(prefix, 'lib/clc')
-
-includedir = options.includedir
-if not includedir:
-  includedir = os.path.join(prefix, 'include')
-
-pkgconfigdir = options.pkgconfigdir
-if not pkgconfigdir:
-  pkgconfigdir = os.path.join(prefix, 'share/pkgconfig')
-
-def llvm_config(args):
-  try:
-    proc = Popen([llvm_config_exe] + args, stdout=PIPE)
-    return proc.communicate()[0].rstrip().replace('\n', ' ')
-  except OSError:
-    print "Error executing llvm-config."
-    print "Please ensure that llvm-config is in your $PATH, or use --with-llvm-config."
-    sys.exit(1)
-
-llvm_version = string.split(string.replace(llvm_config(['--version']), 'svn', ''), '.')
-llvm_int_version = int(llvm_version[0]) * 100 + int(llvm_version[1]) * 10
-llvm_string_version = 'LLVM' + llvm_version[0] + '.' + llvm_version[1]
-
-if llvm_int_version < 400:
-    print "libclc requires LLVM >= 4.0"
-    sys.exit(1)
-
-llvm_system_libs = llvm_config(['--system-libs'])
-llvm_bindir = llvm_config(['--bindir'])
-llvm_core_libs = llvm_config(['--libs', 'core', 'bitreader', 'bitwriter']) + ' ' + \
-                 llvm_system_libs + ' ' + \
-                 llvm_config(['--ldflags'])
-llvm_cxxflags = llvm_config(['--cxxflags']) + ' -fno-exceptions -fno-rtti'
-llvm_libdir = llvm_config(['--libdir'])
-
-llvm_clang = os.path.join(llvm_bindir, 'clang')
-llvm_link = os.path.join(llvm_bindir, 'llvm-link')
-llvm_opt = os.path.join(llvm_bindir, 'opt')
-
-cxx_compiler = options.with_cxx_compiler
-if not cxx_compiler:
-  cxx_compiler = os.path.join(llvm_bindir, 'clang++')
-
-available_targets = {
-  'r600--' : { 'devices' :
-               [{'gpu' : 'cedar',   'aliases' : ['palm', 'sumo', 'sumo2', 'redwood', 'juniper']},
-                {'gpu' : 'cypress', 'aliases' : ['hemlock'] },
-                {'gpu' : 'barts',   'aliases' : ['turks', 'caicos'] },
-                {'gpu' : 'cayman',  'aliases' : ['aruba']} ]},
-  'amdgcn--': { 'devices' :
-                [{'gpu' : 'tahiti', 'aliases' : ['pitcairn', 'verde', 'oland', 'hainan', 'bonaire', 'kabini', 'kaveri', 'hawaii','mullins','tonga','carrizo','iceland','fiji','stoney','polaris10','polaris11']} ]},
-  'amdgcn--amdhsa': { 'devices' :
-                      [{'gpu' : '', 'aliases' : ['bonaire', 'hawaii', 'kabini', 'kaveri', 'mullins', 'carrizo', 'stoney', 'fiji', 'iceland', 'tonga','polaris10','polaris11']} ]},
-  'nvptx--'   : { 'devices' : [{'gpu' : '', 'aliases' : []} ]},
-  'nvptx64--' : { 'devices' : [{'gpu' : '', 'aliases' : []} ]},
-  'nvptx--nvidiacl'   : { 'devices' : [{'gpu' : '', 'aliases' : []} ]},
-  'nvptx64--nvidiacl' : { 'devices' : [{'gpu' : '', 'aliases' : []} ]},
-}
-
-available_targets['amdgcn-mesa-mesa3d'] = available_targets['amdgcn--']
-
-default_targets = ['nvptx--nvidiacl', 'nvptx64--nvidiacl', 'r600--', 'amdgcn--', 'amdgcn--amdhsa', 'amdgcn-mesa-mesa3d']
-
-targets = args
-if not targets:
-  targets = default_targets
-
-b = metabuild.from_name(options.g)
-
-b.rule("LLVM_AS", "%s -o $out $in" % os.path.join(llvm_bindir, "llvm-as"),
-       'LLVM-AS $out')
-b.rule("LLVM_LINK", command = llvm_link + " -o $out $in",
-       description = 'LLVM-LINK $out')
-b.rule("OPT", command = llvm_opt + " -O3 -o $out $in",
-       description = 'OPT $out')
-
-c_compiler_rule(b, "LLVM_TOOL_CXX", 'CXX', cxx_compiler, llvm_cxxflags)
-b.rule("LLVM_TOOL_LINK", cxx_compiler + " -o $out $in %s" % llvm_core_libs + " -Wl,-rpath %s" % llvm_libdir, 'LINK $out')
-
-prepare_builtins = os.path.join('utils', 'prepare-builtins')
-b.build(os.path.join('utils', 'prepare-builtins.o'), "LLVM_TOOL_CXX",
-        os.path.join(srcdir, 'utils', 'prepare-builtins.cpp'))
-b.build(prepare_builtins, "LLVM_TOOL_LINK",
-        os.path.join('utils', 'prepare-builtins.o'))
-
-b.rule("PREPARE_BUILTINS", "%s -o $out $in" % prepare_builtins,
-       'PREPARE-BUILTINS $out')
-b.rule("PYTHON_GEN", "python < $in > $out", "PYTHON_GEN $out")
-b.build('generic/lib/convert.cl', "PYTHON_GEN", ['generic/lib/gen_convert.py'])
-
-manifest_deps = set([sys.argv[0], os.path.join(srcdir, 'build', 'metabuild.py'),
-                     os.path.join(srcdir, 'build', 'ninja_syntax.py')])
-
-install_files_bc = []
-install_deps = []
-
-# Create rules for subnormal helper objects
-for src in ['subnormal_disable.ll', 'subnormal_use_default.ll']:
-  obj_name = src[:-2] + 'bc'
-  obj = os.path.join('generic--', 'lib', obj_name)
-  src_file = os.path.join('generic', 'lib', src)
-  b.build(obj, 'LLVM_AS', src_file)
-  b.default(obj)
-  install_files_bc.append((obj, obj))
-  install_deps.append(obj)
-
-# Create libclc.pc
-clc = open('libclc.pc', 'w')
-clc.write('includedir=%(inc)s\nlibexecdir=%(lib)s\n\nName: libclc\nDescription: Library requirements of the OpenCL C programming language\nVersion: %(maj)s.%(min)s.%(pat)s\nCflags: -I${includedir}\nLibs: -L${libexecdir}' %
-{'inc': includedir, 'lib': libexecdir, 'maj': version_major, 'min': version_minor, 'pat': version_patch})
-clc.close()
-
-for target in targets:
-  (t_arch, t_vendor, t_os) = target.split('-')
-  archs = [t_arch]
-  if t_arch == 'nvptx' or t_arch == 'nvptx64':
-    archs.append('ptx')
-  archs.append('generic')
-
-  subdirs = []
-  for arch in archs:
-    subdirs.append("%s-%s-%s" % (arch, t_vendor, t_os))
-    subdirs.append("%s-%s" % (arch, t_os))
-    if t_os == 'mesa3d':
-        subdirs.append('amdgcn-amdhsa')
-    subdirs.append(arch)
-    if arch == 'amdgcn' or arch == 'r600':
-        subdirs.append('amdgpu')
-
-  incdirs = filter(os.path.isdir,
-               [os.path.join(srcdir, subdir, 'include') for subdir in subdirs])
-  libdirs = filter(lambda d: os.path.isfile(os.path.join(d, 'SOURCES')),
-                   [os.path.join(srcdir, subdir, 'lib') for subdir in subdirs])
-
-  clang_cl_includes = ' '.join(["-I%s" % incdir for incdir in incdirs])
-
-  for device in available_targets[target]['devices']:
-    # The rule for building a .bc file for the specified architecture using clang.
-    clang_bc_flags = "-target %s -I`dirname $in` %s " \
-                     "-fno-builtin " \
-                     "-D__CLC_INTERNAL " \
-                     "-emit-llvm" % (target, clang_cl_includes)
-    if device['gpu'] != '':
-      clang_bc_flags += ' -mcpu=' + device['gpu']
-    clang_bc_rule = "CLANG_CL_BC_" + target + "_" + device['gpu']
-    c_compiler_rule(b, clang_bc_rule, "LLVM-CC", llvm_clang, clang_bc_flags)
-
-    objects = []
-    sources_seen = set()
-    compats_seen = set()
-
-    if device['gpu'] == '':
-      full_target_name = target
-      obj_suffix = ''
-    else:
-      full_target_name = device['gpu'] + '-' + target
-      obj_suffix = '.' + device['gpu']
-
-    for libdir in libdirs:
-      subdir_list_file = os.path.join(libdir, 'SOURCES')
-      manifest_deps.add(subdir_list_file)
-      override_list_file = os.path.join(libdir, 'OVERRIDES')
-      compat_list_file = os.path.join(libdir,
-        'SOURCES_' + llvm_string_version)
-
-      # Build compat list
-      if os.path.exists(compat_list_file):
-        for compat in open(compat_list_file).readlines():
-          compat = compat.rstrip()
-          compats_seen.add(compat)
-
-      # Add target overrides
-      if os.path.exists(override_list_file):
-        for override in open(override_list_file).readlines():
-          override = override.rstrip()
-          sources_seen.add(override)
-
-      for src in open(subdir_list_file).readlines():
-        src = src.rstrip()
-        if src not in sources_seen:
-          sources_seen.add(src)
-          obj = os.path.join(target, 'lib', src + obj_suffix + '.bc')
-          objects.append(obj)
-          src_path = libdir
-          if src in compats_seen:
-            src_path = os.path.join(libdir, llvm_string_version)
-          src_file = os.path.join(src_path, src)
-          ext = os.path.splitext(src)[1]
-          if ext == '.ll':
-            b.build(obj, 'LLVM_AS', src_file)
-          else:
-            b.build(obj, clang_bc_rule, src_file)
-
-    obj = os.path.join('generic--', 'lib', 'subnormal_use_default.bc')
-    if  not options.enable_runtime_subnormal:
-      objects.append(obj)
-
-    builtins_link_bc = os.path.join(target, 'lib', 'builtins.link' + obj_suffix + '.bc')
-    builtins_opt_bc = os.path.join(target, 'lib', 'builtins.opt' + obj_suffix + '.bc')
-    builtins_bc = os.path.join('built_libs', full_target_name + '.bc')
-    b.build(builtins_link_bc, "LLVM_LINK", objects)
-    b.build(builtins_opt_bc, "OPT", builtins_link_bc)
-    b.build(builtins_bc, "PREPARE_BUILTINS", builtins_opt_bc, prepare_builtins)
-    install_files_bc.append((builtins_bc, builtins_bc))
-    install_deps.append(builtins_bc)
-    for alias in device['aliases']:
-      # Ninja cannot have multiple rules with same name so append suffix
-      ruleName = "CREATE_ALIAS_{0}_for_{1}".format(alias, device['gpu'])
-      b.rule(ruleName, "ln -fs %s $out" % os.path.basename(builtins_bc)
-             ,"CREATE-ALIAS $out")
-
-      alias_file = os.path.join('built_libs', alias + '-' + target + '.bc')
-      b.build(alias_file, ruleName, builtins_bc)
-      install_files_bc.append((alias_file, alias_file))
-      install_deps.append(alias_file)
-    b.default(builtins_bc)
-
-
-install_cmd = ' && '.join(['mkdir -p ${DESTDIR}/%(dst)s && cp -r %(src)s ${DESTDIR}/%(dst)s' %
-                           {'src': file,
-                            'dst': libexecdir}
-                           for (file, dest) in install_files_bc])
-install_cmd = ' && '.join(['%(old)s && mkdir -p ${DESTDIR}/%(dst)s && cp -r %(srcdir)s/generic/include/clc ${DESTDIR}/%(dst)s' %
-                           {'old': install_cmd,
-                            'dst': includedir,
-                            'srcdir': srcdir}])
-install_cmd = ' && '.join(['%(old)s && mkdir -p ${DESTDIR}/%(dst)s && cp -r libclc.pc ${DESTDIR}/%(dst)s' %
-                           {'old': install_cmd, 
-                            'dst': pkgconfigdir}])
-  
-b.rule('install', command = install_cmd, description = 'INSTALL')
-b.build('install', 'install', install_deps)
-
-b.rule("configure", command = ' '.join(sys.argv), description = 'CONFIGURE',
-       generator = True)
-b.build(b.output_filename(), 'configure', list(manifest_deps))
-
-b.finish()

libclc/generic/include/clc/as_type.h

@@ -1,68 +0,0 @@
-#define as_char(x) __builtin_astype(x, char)
-#define as_uchar(x) __builtin_astype(x, uchar)
-#define as_short(x) __builtin_astype(x, short)
-#define as_ushort(x) __builtin_astype(x, ushort)
-#define as_int(x) __builtin_astype(x, int)
-#define as_uint(x) __builtin_astype(x, uint)
-#define as_long(x) __builtin_astype(x, long)
-#define as_ulong(x) __builtin_astype(x, ulong)
-#define as_float(x) __builtin_astype(x, float)
-
-#define as_char2(x) __builtin_astype(x, char2)
-#define as_uchar2(x) __builtin_astype(x, uchar2)
-#define as_short2(x) __builtin_astype(x, short2)
-#define as_ushort2(x) __builtin_astype(x, ushort2)
-#define as_int2(x) __builtin_astype(x, int2)
-#define as_uint2(x) __builtin_astype(x, uint2)
-#define as_long2(x) __builtin_astype(x, long2)
-#define as_ulong2(x) __builtin_astype(x, ulong2)
-#define as_float2(x) __builtin_astype(x, float2)
-
-#define as_char3(x) __builtin_astype(x, char3)
-#define as_uchar3(x) __builtin_astype(x, uchar3)
-#define as_short3(x) __builtin_astype(x, short3)
-#define as_ushort3(x) __builtin_astype(x, ushort3)
-#define as_int3(x) __builtin_astype(x, int3)
-#define as_uint3(x) __builtin_astype(x, uint3)
-#define as_long3(x) __builtin_astype(x, long3)
-#define as_ulong3(x) __builtin_astype(x, ulong3)
-#define as_float3(x) __builtin_astype(x, float3)
-
-#define as_char4(x) __builtin_astype(x, char4)
-#define as_uchar4(x) __builtin_astype(x, uchar4)
-#define as_short4(x) __builtin_astype(x, short4)
-#define as_ushort4(x) __builtin_astype(x, ushort4)
-#define as_int4(x) __builtin_astype(x, int4)
-#define as_uint4(x) __builtin_astype(x, uint4)
-#define as_long4(x) __builtin_astype(x, long4)
-#define as_ulong4(x) __builtin_astype(x, ulong4)
-#define as_float4(x) __builtin_astype(x, float4)
-
-#define as_char8(x) __builtin_astype(x, char8)
-#define as_uchar8(x) __builtin_astype(x, uchar8)
-#define as_short8(x) __builtin_astype(x, short8)
-#define as_ushort8(x) __builtin_astype(x, ushort8)
-#define as_int8(x) __builtin_astype(x, int8)
-#define as_uint8(x) __builtin_astype(x, uint8)
-#define as_long8(x) __builtin_astype(x, long8)
-#define as_ulong8(x) __builtin_astype(x, ulong8)
-#define as_float8(x) __builtin_astype(x, float8)
-
-#define as_char16(x) __builtin_astype(x, char16)
-#define as_uchar16(x) __builtin_astype(x, uchar16)
-#define as_short16(x) __builtin_astype(x, short16)
-#define as_ushort16(x) __builtin_astype(x, ushort16)
-#define as_int16(x) __builtin_astype(x, int16)
-#define as_uint16(x) __builtin_astype(x, uint16)
-#define as_long16(x) __builtin_astype(x, long16)
-#define as_ulong16(x) __builtin_astype(x, ulong16)
-#define as_float16(x) __builtin_astype(x, float16)
-
-#ifdef cl_khr_fp64
-#define as_double(x) __builtin_astype(x, double)
-#define as_double2(x) __builtin_astype(x, double2)
-#define as_double3(x) __builtin_astype(x, double3)
-#define as_double4(x) __builtin_astype(x, double4)
-#define as_double8(x) __builtin_astype(x, double8)
-#define as_double16(x) __builtin_astype(x, double16)
-#endif

libclc/generic/include/clc/async/async_work_group_copy.h

@@ -1,15 +0,0 @@
-#define __CLC_DST_ADDR_SPACE local
-#define __CLC_SRC_ADDR_SPACE global
-#define __CLC_BODY <clc/async/async_work_group_copy.inc>
-#include <clc/async/gentype.inc>
-#undef __CLC_DST_ADDR_SPACE
-#undef __CLC_SRC_ADDR_SPACE
-#undef __CLC_BODY
-
-#define __CLC_DST_ADDR_SPACE global
-#define __CLC_SRC_ADDR_SPACE local
-#define __CLC_BODY <clc/async/async_work_group_copy.inc>
-#include <clc/async/gentype.inc>
-#undef __CLC_DST_ADDR_SPACE
-#undef __CLC_SRC_ADDR_SPACE
-#undef __CLC_BODY

libclc/generic/include/clc/async/async_work_group_copy.inc

@@ -1,5 +0,0 @@
-_CLC_OVERLOAD _CLC_DECL event_t async_work_group_copy(
-  __CLC_DST_ADDR_SPACE __CLC_GENTYPE *dst,
-  const __CLC_SRC_ADDR_SPACE __CLC_GENTYPE *src,
-  size_t num_gentypes,
-  event_t event);

libclc/generic/include/clc/async/async_work_group_strided_copy.h

@@ -1,15 +0,0 @@
-#define __CLC_DST_ADDR_SPACE local
-#define __CLC_SRC_ADDR_SPACE global
-#define __CLC_BODY <clc/async/async_work_group_strided_copy.inc>
-#include <clc/async/gentype.inc>
-#undef __CLC_DST_ADDR_SPACE
-#undef __CLC_SRC_ADDR_SPACE
-#undef __CLC_BODY
-
-#define __CLC_DST_ADDR_SPACE global
-#define __CLC_SRC_ADDR_SPACE local
-#define __CLC_BODY <clc/async/async_work_group_strided_copy.inc>
-#include <clc/async/gentype.inc>
-#undef __CLC_DST_ADDR_SPACE
-#undef __CLC_SRC_ADDR_SPACE
-#undef __CLC_BODY

libclc/generic/include/clc/async/async_work_group_strided_copy.inc

@@ -1,6 +0,0 @@
-_CLC_OVERLOAD _CLC_DECL event_t async_work_group_strided_copy(
-  __CLC_DST_ADDR_SPACE __CLC_GENTYPE *dst,
-  const __CLC_SRC_ADDR_SPACE __CLC_GENTYPE *src,
-  size_t num_gentypes,
-  size_t stride,
-  event_t event);

libclc/generic/include/clc/async/gentype.inc

@@ -1,204 +0,0 @@
-
-#define __CLC_GENTYPE char
-#include __CLC_BODY
-#undef __CLC_GENTYPE
-
-#define __CLC_GENTYPE char2
-#include __CLC_BODY
-#undef __CLC_GENTYPE
-
-#define __CLC_GENTYPE char4
-#include __CLC_BODY
-#undef __CLC_GENTYPE
-
-#define __CLC_GENTYPE char8
-#include __CLC_BODY
-#undef __CLC_GENTYPE
-
-#define __CLC_GENTYPE char16
-#include __CLC_BODY
-#undef __CLC_GENTYPE
-
-#define __CLC_GENTYPE uchar
-#include __CLC_BODY
-#undef __CLC_GENTYPE
-
-#define __CLC_GENTYPE uchar2
-#include __CLC_BODY
-#undef __CLC_GENTYPE
-
-#define __CLC_GENTYPE uchar4
-#include __CLC_BODY
-#undef __CLC_GENTYPE
-
-#define __CLC_GENTYPE uchar8
-#include __CLC_BODY
-#undef __CLC_GENTYPE
-
-#define __CLC_GENTYPE uchar16
-#include __CLC_BODY
-#undef __CLC_GENTYPE
-
-#define __CLC_GENTYPE short
-#include __CLC_BODY
-#undef __CLC_GENTYPE
-
-#define __CLC_GENTYPE short2
-#include __CLC_BODY
-#undef __CLC_GENTYPE
-
-#define __CLC_GENTYPE short4
-#include __CLC_BODY
-#undef __CLC_GENTYPE
-
-#define __CLC_GENTYPE short8
-#include __CLC_BODY
-#undef __CLC_GENTYPE
-
-#define __CLC_GENTYPE short16
-#include __CLC_BODY
-#undef __CLC_GENTYPE
-
-#define __CLC_GENTYPE ushort
-#include __CLC_BODY
-#undef __CLC_GENTYPE
-
-#define __CLC_GENTYPE ushort2
-#include __CLC_BODY
-#undef __CLC_GENTYPE
-
-#define __CLC_GENTYPE ushort4
-#include __CLC_BODY
-#undef __CLC_GENTYPE
-
-#define __CLC_GENTYPE ushort8
-#include __CLC_BODY
-#undef __CLC_GENTYPE
-
-#define __CLC_GENTYPE ushort16
-#include __CLC_BODY
-#undef __CLC_GENTYPE
-
-#define __CLC_GENTYPE int
-#include __CLC_BODY
-#undef __CLC_GENTYPE
-
-#define __CLC_GENTYPE int2
-#include __CLC_BODY
-#undef __CLC_GENTYPE
-
-#define __CLC_GENTYPE int4
-#include __CLC_BODY
-#undef __CLC_GENTYPE
-
-#define __CLC_GENTYPE int8
-#include __CLC_BODY
-#undef __CLC_GENTYPE
-
-#define __CLC_GENTYPE int16
-#include __CLC_BODY
-#undef __CLC_GENTYPE
-
-#define __CLC_GENTYPE uint
-#include __CLC_BODY
-#undef __CLC_GENTYPE
-
-#define __CLC_GENTYPE uint2
-#include __CLC_BODY
-#undef __CLC_GENTYPE
-
-#define __CLC_GENTYPE uint4
-#include __CLC_BODY
-#undef __CLC_GENTYPE
-
-#define __CLC_GENTYPE uint8
-#include __CLC_BODY
-#undef __CLC_GENTYPE
-
-#define __CLC_GENTYPE uint16
-#include __CLC_BODY
-#undef __CLC_GENTYPE
-
-#define __CLC_GENTYPE float
-#include __CLC_BODY
-#undef __CLC_GENTYPE
-
-#define __CLC_GENTYPE float2
-#include __CLC_BODY
-#undef __CLC_GENTYPE
-
-#define __CLC_GENTYPE float4
-#include __CLC_BODY
-#undef __CLC_GENTYPE
-
-#define __CLC_GENTYPE float8
-#include __CLC_BODY
-#undef __CLC_GENTYPE
-
-#define __CLC_GENTYPE float16
-#include __CLC_BODY
-#undef __CLC_GENTYPE
-
-#define __CLC_GENTYPE long
-#include __CLC_BODY
-#undef __CLC_GENTYPE
-
-#define __CLC_GENTYPE long2
-#include __CLC_BODY
-#undef __CLC_GENTYPE
-
-#define __CLC_GENTYPE long4
-#include __CLC_BODY
-#undef __CLC_GENTYPE
-
-#define __CLC_GENTYPE long8
-#include __CLC_BODY
-#undef __CLC_GENTYPE
-
-#define __CLC_GENTYPE long16
-#include __CLC_BODY
-#undef __CLC_GENTYPE
-
-#define __CLC_GENTYPE ulong
-#include __CLC_BODY
-#undef __CLC_GENTYPE
-
-#define __CLC_GENTYPE ulong2
-#include __CLC_BODY
-#undef __CLC_GENTYPE
-
-#define __CLC_GENTYPE ulong4
-#include __CLC_BODY
-#undef __CLC_GENTYPE
-
-#define __CLC_GENTYPE ulong8
-#include __CLC_BODY
-#undef __CLC_GENTYPE
-
-#define __CLC_GENTYPE ulong16
-#include __CLC_BODY
-#undef __CLC_GENTYPE
-
-#ifdef cl_khr_fp64
-
-#define __CLC_GENTYPE double
-#include __CLC_BODY
-#undef __CLC_GENTYPE
-
-#define __CLC_GENTYPE double2
-#include __CLC_BODY
-#undef __CLC_GENTYPE
-
-#define __CLC_GENTYPE double4
-#include __CLC_BODY
-#undef __CLC_GENTYPE
-
-#define __CLC_GENTYPE double8
-#include __CLC_BODY
-#undef __CLC_GENTYPE
-
-#define __CLC_GENTYPE double16
-#include __CLC_BODY
-#undef __CLC_GENTYPE
-
-#endif

libclc/generic/include/clc/async/prefetch.h

@@ -1,3 +0,0 @@
-#define __CLC_BODY <clc/async/prefetch.inc>
-#include <clc/async/gentype.inc>
-#undef __CLC_BODY

libclc/generic/include/clc/async/prefetch.inc

@@ -1 +0,0 @@
-_CLC_OVERLOAD _CLC_DECL void prefetch(const global __CLC_GENTYPE *p, size_t num_gentypes);