Port r292146.

Give priority to linker scripts over preemption.

LLD exports symbols that are also present in used shared libraries to
make sure they are preempted at runtime. That is a reasonable default,
but we must allow for it to be overwritten with linker script. If we
don't, libraries that expect to be able to hide a c++ delete operator
will fail.

This should fix the firebird build.

llvm-svn: 292370

clang-tools-extra/docs/ReleaseNotes.rst

@@ -27,11 +27,6 @@ For more information about Clang or LLVM, including information about
 the latest release, please see the `Clang Web Site <http://clang.llvm.org>`_ or
 the `LLVM Web Site <http://llvm.org>`_.
 
-Note that if you are reading this file from a Subversion checkout or the
-main Clang web page, this document applies to the *next* release, not
-the current one. To see the release notes for a specific release, please
-see the `releases page <http://llvm.org/releases/>`_.
-
 What's New in Extra Clang Tools 4.0.0?
 ======================================
 

clang/bindings/python/clang/__init__.py

@@ -20,5 +20,13 @@ The available modules are:
     Bindings for the Clang indexing library.
 """
 
+
+# Python 3 uses unicode for strings. The bindings, in particular the interaction
+# with ctypes, need modifying to handle conversions between unicode and
+# c-strings.
+import sys 
+if sys.version_info[0] != 2: 
+    raise Exception("Only Python 2 is supported.")
+
 __all__ = ['cindex']
 

clang/bindings/python/clang/cindex.py

@@ -64,7 +64,6 @@ call is efficient.
 
 from ctypes import *
 import collections
-import sys
 
 import clang.enumerations
 
@@ -74,33 +73,6 @@ import clang.enumerations
 # this by marshalling object arguments as void**.
 c_object_p = POINTER(c_void_p)
 
-if sys.version_info[0] > 2:
-# Python 3 strings are unicode, translate them to/from utf8 for C-interop
-# Python 3 replaces xrange with range, we want xrange behaviour
-    xrange = range
-
-    class c_string_p(c_char_p):
-        def __init__(self, p=None):
-            if type(p) == str:
-                p = p.encode("utf8")
-            super(c_char_p, self).__init__(p)
-
-        def __str__(self):
-            return str(self.value)
-
-        @property
-        def value(self):
-            if super(c_char_p, self).value is None:
-                return None
-            return super(c_char_p, self).value.decode("utf8")
-
-        @classmethod
-        def from_param(cls, param):
-                return cls(param)
-else:
-    c_string_p = c_char_p
-
-
 callbacks = {}
 
 ### Exception Classes ###
@@ -175,7 +147,7 @@ class CachedProperty(object):
 class _CXString(Structure):
     """Helper for transforming CXString results."""
 
-    _fields_ = [("spelling", c_string_p), ("free", c_int)]
+    _fields_ = [("spelling", c_char_p), ("free", c_int)]
 
     def __del__(self):
         conf.lib.clang_disposeString(self)
@@ -357,7 +329,7 @@ class Diagnostic(object):
 
     @property
     def spelling(self):
-        return str(conf.lib.clang_getDiagnosticSpelling(self))
+        return conf.lib.clang_getDiagnosticSpelling(self)
 
     @property
     def ranges(self):
@@ -386,8 +358,8 @@ class Diagnostic(object):
 
             def __getitem__(self, key):
                 range = SourceRange()
-                value = str(conf.lib.clang_getDiagnosticFixIt(self.diag, key,
+                value = conf.lib.clang_getDiagnosticFixIt(self.diag, key,
-                        byref(range)))
+                        byref(range))
                 if len(value) == 0:
                     raise IndexError
 
@@ -420,12 +392,12 @@ class Diagnostic(object):
     @property
     def category_name(self):
         """The string name of the category for this diagnostic."""
-        return str(conf.lib.clang_getDiagnosticCategoryText(self))
+        return conf.lib.clang_getDiagnosticCategoryText(self)
 
     @property
     def option(self):
         """The command-line option that enables this diagnostic."""
-        return str(conf.lib.clang_getDiagnosticOption(self, None))
+        return conf.lib.clang_getDiagnosticOption(self, None)
 
     @property
     def disable_option(self):
@@ -433,7 +405,7 @@ class Diagnostic(object):
         disable = _CXString()
         conf.lib.clang_getDiagnosticOption(self, byref(disable))
 
-        return str(conf.lib.clang_getCString(disable))
+        return conf.lib.clang_getCString(disable)
 
     def format(self, options=None):
         """
@@ -600,7 +572,7 @@ class BaseEnumeration(object):
             for key, value in self.__class__.__dict__.items():
                 if isinstance(value, self.__class__):
                     self._name_map[value] = key
-        return str(self._name_map[self])
+        return self._name_map[self]
 
     @classmethod
     def from_id(cls, id):
@@ -624,7 +596,7 @@ class CursorKind(BaseEnumeration):
     @staticmethod
     def get_all_kinds():
         """Return all CursorKind enumeration instances."""
-        return [x for x in CursorKind._kinds if x]
+        return filter(None, CursorKind._kinds)
 
     def is_declaration(self):
         """Test if this is a declaration kind."""
@@ -1457,9 +1429,9 @@ class Cursor(Structure):
     def spelling(self):
         """Return the spelling of the entity pointed at by the cursor."""
         if not hasattr(self, '_spelling'):
-            self._spelling = str(conf.lib.clang_getCursorSpelling(self))
+            self._spelling = conf.lib.clang_getCursorSpelling(self)
 
-        return str(self._spelling)
+        return self._spelling
 
     @property
     def displayname(self):
@@ -1471,7 +1443,7 @@ class Cursor(Structure):
         arguments of a class template specialization.
         """
         if not hasattr(self, '_displayname'):
-            self._displayname = str(conf.lib.clang_getCursorDisplayName(self))
+            self._displayname = conf.lib.clang_getCursorDisplayName(self)
 
         return self._displayname
 
@@ -1479,7 +1451,7 @@ class Cursor(Structure):
     def mangled_name(self):
         """Return the mangled name for the entity referenced by this cursor."""
         if not hasattr(self, '_mangled_name'):
-            self._mangled_name = str(conf.lib.clang_Cursor_getMangling(self))
+            self._mangled_name = conf.lib.clang_Cursor_getMangling(self)
 
         return self._mangled_name
 
@@ -1618,7 +1590,7 @@ class Cursor(Structure):
             self._objc_type_encoding = \
               conf.lib.clang_getDeclObjCTypeEncoding(self)
 
-        return str(self._objc_type_encoding)
+        return self._objc_type_encoding
 
     @property
     def hash(self):
@@ -1665,23 +1637,17 @@ class Cursor(Structure):
     @property
     def brief_comment(self):
         """Returns the brief comment text associated with that Cursor"""
-        r = conf.lib.clang_Cursor_getBriefCommentText(self)
+        return conf.lib.clang_Cursor_getBriefCommentText(self)
-        if not r:
-            return None    
-        return str(r)
 
     @property
     def raw_comment(self):
         """Returns the raw comment text associated with that Cursor"""
-        r = conf.lib.clang_Cursor_getRawCommentText(self)
+        return conf.lib.clang_Cursor_getRawCommentText(self)
-        if not r:
-            return None
-        return str(r)
 
     def get_arguments(self):
         """Return an iterator for accessing the arguments of this cursor."""
         num_args = conf.lib.clang_Cursor_getNumArguments(self)
-        for i in xrange(0, num_args):
+        for i in range(0, num_args):
             yield conf.lib.clang_Cursor_getArgument(self, i)
 
     def get_num_template_arguments(self):
@@ -1885,7 +1851,7 @@ class TypeKind(BaseEnumeration):
     @property
     def spelling(self):
         """Retrieve the spelling of this TypeKind."""
-        return str(conf.lib.clang_getTypeKindSpelling(self.value))
+        return conf.lib.clang_getTypeKindSpelling(self.value)
 
     def __repr__(self):
         return 'TypeKind.%s' % (self.name,)
@@ -2161,7 +2127,7 @@ class Type(Structure):
         """
         Retrieve the offset of a field in the record.
         """
-        return conf.lib.clang_Type_getOffsetOf(self, fieldname)
+        return conf.lib.clang_Type_getOffsetOf(self, c_char_p(fieldname))
 
     def get_ref_qualifier(self):
         """
@@ -2188,7 +2154,7 @@ class Type(Structure):
     @property
     def spelling(self):
         """Retrieve the spelling of this Type."""
-        return str(conf.lib.clang_getTypeSpelling(self))
+        return conf.lib.clang_getTypeSpelling(self)
 
     def __eq__(self, other):
         if type(other) != type(self):
@@ -2220,7 +2186,7 @@ class ClangObject(object):
 
 class _CXUnsavedFile(Structure):
     """Helper for passing unsaved file arguments."""
-    _fields_ = [("name", c_string_p), ("contents", c_string_p), ('length', c_ulong)]
+    _fields_ = [("name", c_char_p), ("contents", c_char_p), ('length', c_ulong)]
 
 # Functions calls through the python interface are rather slow. Fortunately,
 # for most symboles, we do not need to perform a function call. Their spelling
@@ -2266,7 +2232,7 @@ class CompletionChunk:
         self.__kindNumberCache = -1
 
     def __repr__(self):
-        return "{'" + str(self.spelling) + "', " + str(self.kind) + "}"
+        return "{'" + self.spelling + "', " + str(self.kind) + "}"
 
     @CachedProperty
     def spelling(self):
@@ -2575,9 +2541,7 @@ class TranslationUnit(ClangObject):
 
         args_array = None
         if len(args) > 0:
-            args_array = (c_string_p * len(args))()
+            args_array = (c_char_p * len(args))(* args)
-            for i,a in enumerate(args):
-                args_array[i] = c_string_p(a)
 
         unsaved_array = None
         if len(unsaved_files) > 0:
@@ -2586,8 +2550,8 @@ class TranslationUnit(ClangObject):
                 if hasattr(contents, "read"):
                     contents = contents.read()
 
-                unsaved_array[i].name = c_string_p(name)
+                unsaved_array[i].name = name
-                unsaved_array[i].contents = c_string_p(contents)
+                unsaved_array[i].contents = contents
                 unsaved_array[i].length = len(contents)
 
         ptr = conf.lib.clang_parseTranslationUnit(index, filename, args_array,
@@ -2642,7 +2606,7 @@ class TranslationUnit(ClangObject):
     @property
     def spelling(self):
         """Get the original translation unit source file name."""
-        return str(conf.lib.clang_getTranslationUnitSpelling(self))
+        return conf.lib.clang_getTranslationUnitSpelling(self)
 
     def get_includes(self):
         """
@@ -2832,8 +2796,8 @@ class TranslationUnit(ClangObject):
                     print(value)
                 if not isinstance(value, str):
                     raise TypeError('Unexpected unsaved file contents.')
-                unsaved_files_array[i].name = c_string_p(name)
+                unsaved_files_array[i].name = name
-                unsaved_files_array[i].contents = c_string_p(value)
+                unsaved_files_array[i].contents = value
                 unsaved_files_array[i].length = len(value)
         ptr = conf.lib.clang_codeCompleteAt(self, path, line, column,
                 unsaved_files_array, len(unsaved_files), options)
@@ -2868,7 +2832,7 @@ class File(ClangObject):
     @property
     def name(self):
         """Return the complete file and path name of the file."""
-        return str(conf.lib.clang_getCString(conf.lib.clang_getFileName(self)))
+        return conf.lib.clang_getCString(conf.lib.clang_getFileName(self))
 
     @property
     def time(self):
@@ -2876,7 +2840,7 @@ class File(ClangObject):
         return conf.lib.clang_getFileTime(self)
 
     def __str__(self):
-        return str(self.name)
+        return self.name
 
     def __repr__(self):
         return "<File: %s>" % (self.name)
@@ -2945,12 +2909,12 @@ class CompileCommand(object):
     @property
     def directory(self):
         """Get the working directory for this CompileCommand"""
-        return str(conf.lib.clang_CompileCommand_getDirectory(self.cmd))
+        return conf.lib.clang_CompileCommand_getDirectory(self.cmd)
 
     @property
     def filename(self):
         """Get the working filename for this CompileCommand"""
-        return str(conf.lib.clang_CompileCommand_getFilename(self.cmd))
+        return conf.lib.clang_CompileCommand_getFilename(self.cmd)
 
     @property
     def arguments(self):
@@ -2962,7 +2926,7 @@ class CompileCommand(object):
         """
         length = conf.lib.clang_CompileCommand_getNumArgs(self.cmd)
         for i in xrange(length):
-            yield str(conf.lib.clang_CompileCommand_getArg(self.cmd, i))
+            yield conf.lib.clang_CompileCommand_getArg(self.cmd, i)
 
 class CompileCommands(object):
     """
@@ -3056,7 +3020,7 @@ class Token(Structure):
 
         This is the textual representation of the token in source.
         """
-        return str(conf.lib.clang_getTokenSpelling(self._tu, self))
+        return conf.lib.clang_getTokenSpelling(self._tu, self)
 
     @property
     def kind(self):
@@ -3099,7 +3063,7 @@ functionList = [
    [c_object_p]),
 
   ("clang_CompilationDatabase_fromDirectory",
-   [c_string_p, POINTER(c_uint)],
+   [c_char_p, POINTER(c_uint)],
    c_object_p,
    CompilationDatabase.from_result),
 
@@ -3109,7 +3073,7 @@ functionList = [
    CompileCommands.from_result),
 
   ("clang_CompilationDatabase_getCompileCommands",
-   [c_object_p, c_string_p],
+   [c_object_p, c_char_p],
    c_object_p,
    CompileCommands.from_result),
 
@@ -3144,7 +3108,7 @@ functionList = [
    c_uint),
 
   ("clang_codeCompleteAt",
-   [TranslationUnit, c_string_p, c_int, c_int, c_void_p, c_int, c_int],
+   [TranslationUnit, c_char_p, c_int, c_int, c_void_p, c_int, c_int],
    POINTER(CCRStructure)),
 
   ("clang_codeCompleteGetDiagnostic",
@@ -3160,7 +3124,7 @@ functionList = [
    c_object_p),
 
   ("clang_createTranslationUnit",
-   [Index, c_string_p],
+   [Index, c_char_p],
    c_object_p),
 
   ("clang_CXXConstructor_isConvertingConstructor",
@@ -3310,7 +3274,7 @@ functionList = [
 
   ("clang_getCString",
    [_CXString],
-   c_string_p),
+   c_char_p),
 
   ("clang_getCursor",
    [TranslationUnit, SourceLocation],
@@ -3457,7 +3421,7 @@ functionList = [
    Type.from_result),
 
   ("clang_getFile",
-   [TranslationUnit, c_string_p],
+   [TranslationUnit, c_char_p],
    c_object_p),
 
   ("clang_getFileName",
@@ -3586,7 +3550,7 @@ functionList = [
 
   ("clang_getTUResourceUsageName",
    [c_uint],
-   c_string_p),
+   c_char_p),
 
   ("clang_getTypeDeclaration",
    [Type],
@@ -3681,7 +3645,7 @@ functionList = [
    bool),
 
   ("clang_parseTranslationUnit",
-   [Index, c_string_p, c_void_p, c_int, c_void_p, c_int, c_int],
+   [Index, c_char_p, c_void_p, c_int, c_void_p, c_int, c_int],
    c_object_p),
 
   ("clang_reparseTranslationUnit",
@@ -3689,7 +3653,7 @@ functionList = [
    c_int),
 
   ("clang_saveTranslationUnit",
-   [TranslationUnit, c_string_p, c_uint],
+   [TranslationUnit, c_char_p, c_uint],
    c_int),
 
   ("clang_tokenize",
@@ -3761,7 +3725,7 @@ functionList = [
    Type.from_result),
 
   ("clang_Type_getOffsetOf",
-   [Type, c_string_p],
+   [Type, c_char_p],
    c_longlong),
 
   ("clang_Type_getSizeOf",
@@ -3820,8 +3784,7 @@ def register_functions(lib, ignore_errors):
     def register(item):
         return register_function(lib, item, ignore_errors)
 
-    for f in functionList:
+    map(register, functionList)
-        register(f)
 
 class Config:
     library_path = None

clang/bindings/python/tests/cindex/test_translation_unit.py

@@ -59,13 +59,9 @@ int SOME_DEFINE;
     assert spellings[-1] == 'y'
 
 def test_unsaved_files_2():
-    try:
+    import StringIO
-        from StringIO import StringIO
-    except:
-        from io import StringIO
-
     tu = TranslationUnit.from_source('fake.c', unsaved_files = [
-            ('fake.c', StringIO('int x;'))])
+            ('fake.c', StringIO.StringIO('int x;'))])
     spellings = [c.spelling for c in tu.cursor.get_children()]
     assert spellings[-1] == 'x'
 

clang/docs/ReleaseNotes.rst

@@ -31,11 +31,6 @@ the latest release, please check out the main please see the `Clang Web
 Site <http://clang.llvm.org>`_ or the `LLVM Web
 Site <http://llvm.org>`_.
 
-Note that if you are reading this file from a Subversion checkout or the
-main Clang web page, this document applies to the *next* release, not
-the current one. To see the release notes for a specific release, please
-see the `releases page <http://llvm.org/releases/>`_.
-
 What's New in Clang 4.0.0?
 ==========================
 
@@ -51,8 +46,35 @@ Major New Features
   clang to emit a warning or error if a function call meets one or more
   user-specified conditions.
 
+- Enhanced devirtualization with
+  `-fstrict-vtable-pointers <UsersManual.html#cmdoption-fstrict-vtable-pointers>`_.
+  Clang devirtualizes across different basic blocks, like loops:
+
+  .. code-block:: c++
+
+       struct A {
+           virtual void foo();
+       };
+       void indirect(A &a, int n) {
+           for (int i = 0 ; i < n; i++)
+               a.foo();
+       }
+       void test(int n) {
+           A a;
+           indirect(a, n);
+       }
+
+
 -  ...
 
+Improvements to ThinLTO (-flto=thin)
+------------------------------------
+- Integration with profile data (PGO). When available, profile data enables
+  more accurate function importing decisions, as well as cross-module indirect
+  call promotion.
+- Significant build-time and binary-size improvements when compiling with debug
+  info (-g).
+
 Improvements to Clang's diagnostics
 ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
 

clang/docs/UsersManual.rst

@@ -1097,6 +1097,14 @@ are listed below.
    the behavior of sanitizers in the ``cfi`` group to allow checking
    of cross-DSO virtual and indirect calls.
 
+
+.. option:: -fstrict-vtable-pointers
+
+   Enable optimizations based on the strict rules for overwriting polymorphic
+   C++ objects, i.e. the vptr is invariant during an object's lifetime.
+   This enables better devirtualization. Turned off by default, because it is
+   still experimental.
+
 .. option:: -ffast-math
 
    Enable fast-math mode. This defines the ``__FAST_MATH__`` preprocessor
@@ -2645,7 +2653,7 @@ Execute ``clang-cl /?`` to see a list of supported options:
                               (overridden by LLVM_PROFILE_FILE env var)
       -fprofile-instr-generate
                               Generate instrumented code to collect execution counts into default.profraw file
-                              (overriden by '=' form of option or LLVM_PROFILE_FILE env var)
+                              (overridden by '=' form of option or LLVM_PROFILE_FILE env var)
       -fprofile-instr-use=<value>
                               Use instrumentation data for profile-guided optimization
       -fsanitize-blacklist=<value>

clang/include/clang/Basic/AttrDocs.td

@@ -386,6 +386,7 @@ warnings or errors at compile-time if calls to the attributed function meet
 certain user-defined criteria. For example:
 
 .. code-block:: c
+
   void abs(int a)
     __attribute__((diagnose_if(a >= 0, "Redundant abs call", "warning")));
   void must_abs(int a)

clang/lib/Analysis/CFG.cpp

@@ -2175,19 +2175,15 @@ CFGBlock *CFGBuilder::VisitIfStmt(IfStmt *I) {
   SaveAndRestore<LocalScope::const_iterator> save_scope_pos(ScopePos);
 
   // Create local scope for C++17 if init-stmt if one exists.
-  if (Stmt *Init = I->getInit()) {
+  if (Stmt *Init = I->getInit())
-    LocalScope::const_iterator BeginScopePos = ScopePos;
     addLocalScopeForStmt(Init);
-    addAutomaticObjDtors(ScopePos, BeginScopePos, I);
-  }
 
   // Create local scope for possible condition variable.
   // Store scope position. Add implicit destructor.
-  if (VarDecl *VD = I->getConditionVariable()) {
+  if (VarDecl *VD = I->getConditionVariable())
-    LocalScope::const_iterator BeginScopePos = ScopePos;
     addLocalScopeForVarDecl(VD);
-    addAutomaticObjDtors(ScopePos, BeginScopePos, I);
+
-  }
+  addAutomaticObjDtors(ScopePos, save_scope_pos.get(), I);
 
   // The block we were processing is now finished.  Make it the successor
   // block.
@@ -2256,36 +2252,39 @@ CFGBlock *CFGBuilder::VisitIfStmt(IfStmt *I) {
   // removes infeasible paths from the control-flow graph by having the
   // control-flow transfer of '&&' or '||' go directly into the then/else
   // blocks directly.
-  if (!I->getConditionVariable())
+  BinaryOperator *Cond =
-    if (BinaryOperator *Cond =
+      I->getConditionVariable()
-            dyn_cast<BinaryOperator>(I->getCond()->IgnoreParens()))
+          ? nullptr
-      if (Cond->isLogicalOp())
+          : dyn_cast<BinaryOperator>(I->getCond()->IgnoreParens());
-        return VisitLogicalOperator(Cond, I, ThenBlock, ElseBlock).first;
+  CFGBlock *LastBlock;
+  if (Cond && Cond->isLogicalOp())
+    LastBlock = VisitLogicalOperator(Cond, I, ThenBlock, ElseBlock).first;
+  else {
+    // Now create a new block containing the if statement.
+    Block = createBlock(false);
 
-  // Now create a new block containing the if statement.
+    // Set the terminator of the new block to the If statement.
-  Block = createBlock(false);
+    Block->setTerminator(I);
 
-  // Set the terminator of the new block to the If statement.
+    // See if this is a known constant.
-  Block->setTerminator(I);
+    const TryResult &KnownVal = tryEvaluateBool(I->getCond());
 
-  // See if this is a known constant.
+    // Add the successors.  If we know that specific branches are
-  const TryResult &KnownVal = tryEvaluateBool(I->getCond());
+    // unreachable, inform addSuccessor() of that knowledge.
+    addSuccessor(Block, ThenBlock, /* isReachable = */ !KnownVal.isFalse());
+    addSuccessor(Block, ElseBlock, /* isReachable = */ !KnownVal.isTrue());
 
-  // Add the successors.  If we know that specific branches are
+    // Add the condition as the last statement in the new block.  This may
-  // unreachable, inform addSuccessor() of that knowledge.
+    // create new blocks as the condition may contain control-flow.  Any newly
-  addSuccessor(Block, ThenBlock, /* isReachable = */ !KnownVal.isFalse());
+    // created blocks will be pointed to be "Block".
-  addSuccessor(Block, ElseBlock, /* isReachable = */ !KnownVal.isTrue());
+    LastBlock = addStmt(I->getCond());
 
-  // Add the condition as the last statement in the new block.  This may create
+    // If the IfStmt contains a condition variable, add it and its
-  // new blocks as the condition may contain control-flow.  Any newly created
+    // initializer to the CFG.
-  // blocks will be pointed to be "Block".
+    if (const DeclStmt* DS = I->getConditionVariableDeclStmt()) {
-  CFGBlock *LastBlock = addStmt(I->getCond());
+      autoCreateBlock();
-
+      LastBlock = addStmt(const_cast<DeclStmt *>(DS));
-  // If the IfStmt contains a condition variable, add it and its
+    }
-  // initializer to the CFG.
-  if (const DeclStmt* DS = I->getConditionVariableDeclStmt()) {
-    autoCreateBlock();
-    LastBlock = addStmt(const_cast<DeclStmt *>(DS));
   }
 
   // Finally, if the IfStmt contains a C++17 init-stmt, add it to the CFG.
@@ -3078,19 +3077,15 @@ CFGBlock *CFGBuilder::VisitSwitchStmt(SwitchStmt *Terminator) {
   SaveAndRestore<LocalScope::const_iterator> save_scope_pos(ScopePos);
 
   // Create local scope for C++17 switch init-stmt if one exists.
-  if (Stmt *Init = Terminator->getInit()) {
+  if (Stmt *Init = Terminator->getInit())
-    LocalScope::const_iterator BeginScopePos = ScopePos;
     addLocalScopeForStmt(Init);
-    addAutomaticObjDtors(ScopePos, BeginScopePos, Terminator);
-  }
 
   // Create local scope for possible condition variable.
   // Store scope position. Add implicit destructor.
-  if (VarDecl *VD = Terminator->getConditionVariable()) {
+  if (VarDecl *VD = Terminator->getConditionVariable())
-    LocalScope::const_iterator SwitchBeginScopePos = ScopePos;
     addLocalScopeForVarDecl(VD);
-    addAutomaticObjDtors(ScopePos, SwitchBeginScopePos, Terminator);
+
-  }
+  addAutomaticObjDtors(ScopePos, save_scope_pos.get(), Terminator);
 
   if (Block) {
     if (badCFG)

clang/lib/CodeGen/CGDeclCXX.cpp

@@ -353,9 +353,6 @@ CodeGenModule::EmitCXXGlobalVarDeclInitFunc(const VarDecl *D,
 
   if (D->getTLSKind()) {
     // FIXME: Should we support init_priority for thread_local?
-    // FIXME: Ideally, initialization of instantiated thread_local static data
-    // members of class templates should not trigger initialization of other
-    // entities in the TU.
     // FIXME: We only need to register one __cxa_thread_atexit function for the
     // entire TU.
     CXXThreadLocalInits.push_back(Fn);

clang/lib/CodeGen/ItaniumCXXABI.cpp

@@ -2272,7 +2272,21 @@ void ItaniumCXXABI::EmitThreadLocalInitFuncs(
     ArrayRef<llvm::Function *> CXXThreadLocalInits,
     ArrayRef<const VarDecl *> CXXThreadLocalInitVars) {
   llvm::Function *InitFunc = nullptr;
-  if (!CXXThreadLocalInits.empty()) {
+
+  // Separate initializers into those with ordered (or partially-ordered)
+  // initialization and those with unordered initialization.
+  llvm::SmallVector<llvm::Function *, 8> OrderedInits;
+  llvm::SmallDenseMap<const VarDecl *, llvm::Function *> UnorderedInits;
+  for (unsigned I = 0; I != CXXThreadLocalInits.size(); ++I) {
+    if (isTemplateInstantiation(
+            CXXThreadLocalInitVars[I]->getTemplateSpecializationKind()))
+      UnorderedInits[CXXThreadLocalInitVars[I]->getCanonicalDecl()] =
+          CXXThreadLocalInits[I];
+    else
+      OrderedInits.push_back(CXXThreadLocalInits[I]);
+  }
+
+  if (!OrderedInits.empty()) {
     // Generate a guarded initialization function.
     llvm::FunctionType *FTy =
         llvm::FunctionType::get(CGM.VoidTy, /*isVarArg=*/false);
@@ -2289,24 +2303,28 @@ void ItaniumCXXABI::EmitThreadLocalInitFuncs(
     CharUnits GuardAlign = CharUnits::One();
     Guard->setAlignment(GuardAlign.getQuantity());
 
-    CodeGenFunction(CGM)
+    CodeGenFunction(CGM).GenerateCXXGlobalInitFunc(InitFunc, OrderedInits,
-        .GenerateCXXGlobalInitFunc(InitFunc, CXXThreadLocalInits,
+                                                   Address(Guard, GuardAlign));
-                                   Address(Guard, GuardAlign));
     // On Darwin platforms, use CXX_FAST_TLS calling convention.
     if (CGM.getTarget().getTriple().isOSDarwin()) {
       InitFunc->setCallingConv(llvm::CallingConv::CXX_FAST_TLS);
       InitFunc->addFnAttr(llvm::Attribute::NoUnwind);
     }
   }
+
+  // Emit thread wrappers.
   for (const VarDecl *VD : CXXThreadLocals) {
     llvm::GlobalVariable *Var =
         cast<llvm::GlobalVariable>(CGM.GetGlobalValue(CGM.getMangledName(VD)));
+    llvm::Function *Wrapper = getOrCreateThreadLocalWrapper(VD, Var);
 
     // Some targets require that all access to thread local variables go through
     // the thread wrapper.  This means that we cannot attempt to create a thread
     // wrapper or a thread helper.
-    if (isThreadWrapperReplaceable(VD, CGM) && !VD->hasDefinition())
+    if (isThreadWrapperReplaceable(VD, CGM) && !VD->hasDefinition()) {
+      Wrapper->setLinkage(llvm::Function::ExternalLinkage);
       continue;
+    }
 
     // Mangle the name for the thread_local initialization function.
     SmallString<256> InitFnName;
@@ -2322,18 +2340,21 @@ void ItaniumCXXABI::EmitThreadLocalInitFuncs(
     bool InitIsInitFunc = false;
     if (VD->hasDefinition()) {
       InitIsInitFunc = true;
-      if (InitFunc)
+      llvm::Function *InitFuncToUse = InitFunc;
+      if (isTemplateInstantiation(VD->getTemplateSpecializationKind()))
+        InitFuncToUse = UnorderedInits.lookup(VD->getCanonicalDecl());
+      if (InitFuncToUse)
         Init = llvm::GlobalAlias::create(Var->getLinkage(), InitFnName.str(),
-                                         InitFunc);
+                                         InitFuncToUse);
     } else {
       // Emit a weak global function referring to the initialization function.
       // This function will not exist if the TU defining the thread_local
       // variable in question does not need any dynamic initialization for
       // its thread_local variables.
       llvm::FunctionType *FnTy = llvm::FunctionType::get(CGM.VoidTy, false);
-      Init = llvm::Function::Create(
+      Init = llvm::Function::Create(FnTy,
-          FnTy, llvm::GlobalVariable::ExternalWeakLinkage, InitFnName.str(),
+                                    llvm::GlobalVariable::ExternalWeakLinkage,
-          &CGM.getModule());
+                                    InitFnName.str(), &CGM.getModule());
       const CGFunctionInfo &FI = CGM.getTypes().arrangeNullaryFunction();
       CGM.SetLLVMFunctionAttributes(nullptr, FI, cast<llvm::Function>(Init));
     }
@@ -2341,7 +2362,6 @@ void ItaniumCXXABI::EmitThreadLocalInitFuncs(
     if (Init)
       Init->setVisibility(Var->getVisibility());
 
-    llvm::Function *Wrapper = getOrCreateThreadLocalWrapper(VD, Var);
     llvm::LLVMContext &Context = CGM.getModule().getContext();
     llvm::BasicBlock *Entry = llvm::BasicBlock::Create(Context, "", Wrapper);
     CGBuilderTy Builder(CGM, Entry);

clang/lib/Driver/Tools.cpp

@@ -6431,11 +6431,13 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA,
     A->claim();
 
     // We translate this by hand to the -cc1 argument, since nightly test uses
-    // it and developers have been trained to spell it with -mllvm.
+    // it and developers have been trained to spell it with -mllvm. Both
-    if (StringRef(A->getValue(0)) == "-disable-llvm-passes") {
+    // spellings are now deprecated and should be removed.
-      CmdArgs.push_back("-disable-llvm-passes");
+    if (StringRef(A->getValue(0)) == "-disable-llvm-optzns") {
-    } else
+      CmdArgs.push_back("-disable-llvm-optzns");
+    } else {
       A->render(Args, CmdArgs);
+    }
   }
 
   // With -save-temps, we want to save the unoptimized bitcode output from the

clang/lib/Frontend/DependencyFile.cpp

@@ -447,9 +447,9 @@ void DFGImpl::OutputDependencyFile() {
   // Create phony targets if requested.
   if (PhonyTarget && !Files.empty()) {
     // Skip the first entry, this is always the input file itself.
-    for (StringRef File : Files) {
+    for (auto I = Files.begin() + 1, E = Files.end(); I != E; ++I) {
       OS << '\n';
-      PrintFilename(OS, File, OutputFormat);
+      PrintFilename(OS, *I, OutputFormat);
       OS << ":\n";
     }
   }

clang/lib/Parse/ParseExpr.cpp

@@ -1652,9 +1652,10 @@ Parser::ParsePostfixExpressionSuffix(ExprResult LHS) {
 
       if (Tok.is(tok::code_completion)) {
         // Code completion for a member access expression.
-        Actions.CodeCompleteMemberReferenceExpr(
+        if (Expr *Base = LHS.get())
-            getCurScope(), LHS.get(), OpLoc, OpKind == tok::arrow,
+          Actions.CodeCompleteMemberReferenceExpr(
-            ExprStatementTokLoc == LHS.get()->getLocStart());
+              getCurScope(), Base, OpLoc, OpKind == tok::arrow,
+              ExprStatementTokLoc == Base->getLocStart());
 
         cutOffParsing();
         return ExprError();

clang/lib/Sema/SemaTemplate.cpp

@@ -5127,18 +5127,22 @@ ExprResult Sema::CheckTemplateArgument(NonTypeTemplateParmDecl *Param,
   if (CTAK == CTAK_Deduced &&
       !Context.hasSameType(ParamType.getNonLValueExprType(Context),
                            Arg->getType())) {
-    // C++ [temp.deduct.type]p17: (DR1770)
+    // FIXME: If either type is dependent, we skip the check. This isn't
-    //   If P has a form that contains <i>, and if the type of i differs from
+    // correct, since during deduction we're supposed to have replaced each
-    //   the type of the corresponding template parameter of the template named
+    // template parameter with some unique (non-dependent) placeholder.
-    //   by the enclosing simple-template-id, deduction fails.
+    // FIXME: If the argument type contains 'auto', we carry on and fail the
-    //
+    // type check in order to force specific types to be more specialized than
-    // Note that CTAK will be CTAK_DeducedFromArrayBound if the form was [i]
+    // 'auto'. It's not clear how partial ordering with 'auto' is supposed to
-    // rather than <i>.
+    // work.
-    //
+    if ((ParamType->isDependentType() || Arg->isTypeDependent()) &&
-    // FIXME: We interpret the 'i' here as referring to the expression
+        !Arg->getType()->getContainedAutoType()) {
-    // denoting the non-type template parameter rather than the parameter
+      Converted = TemplateArgument(Arg);
-    // itself, and so strip off references before comparing types. It's
+      return Arg;
-    // not clear how this is supposed to work for references.
+    }
+    // FIXME: This attempts to implement C++ [temp.deduct.type]p17. Per DR1770,
+    // we should actually be checking the type of the template argument in P,
+    // not the type of the template argument deduced from A, against the
+    // template parameter type.
     Diag(StartLoc, diag::err_deduced_non_type_template_arg_type_mismatch)
       << Arg->getType()
       << ParamType.getUnqualifiedType();
@@ -7789,6 +7793,7 @@ Sema::ActOnExplicitInstantiation(Scope *S,
   Specialization->setTemplateKeywordLoc(TemplateLoc);
   Specialization->setBraceRange(SourceRange());
 
+  bool PreviouslyDLLExported = Specialization->hasAttr<DLLExportAttr>();
   if (Attr)
     ProcessDeclAttributeList(S, Specialization, Attr);
 
@@ -7851,8 +7856,9 @@ Sema::ActOnExplicitInstantiation(Scope *S,
 
     // Fix a TSK_ImplicitInstantiation followed by a
     // TSK_ExplicitInstantiationDefinition
-    if (Old_TSK == TSK_ImplicitInstantiation &&
+    bool NewlyDLLExported =
-        Specialization->hasAttr<DLLExportAttr>() &&
+        !PreviouslyDLLExported && Specialization->hasAttr<DLLExportAttr>();
+    if (Old_TSK == TSK_ImplicitInstantiation && NewlyDLLExported &&
         (Context.getTargetInfo().getCXXABI().isMicrosoft() ||
          Context.getTargetInfo().getTriple().isWindowsItaniumEnvironment())) {
       // In the MS ABI, an explicit instantiation definition can add a dll

clang/test/CXX/drs/dr0xx.cpp

@@ -1032,7 +1032,7 @@ namespace dr91 { // dr91: yes
   int k = f(U());
 }
 
-namespace dr92 { // dr92: 4.0 c++17
+namespace dr92 { // dr92: 4 c++17
   void f() throw(int, float); // expected-error 0-1{{ISO C++1z does not allow}} expected-note 0-1{{use 'noexcept}}
   void (*p)() throw(int) = &f; // expected-error 0-1{{ISO C++1z does not allow}} expected-note 0-1{{use 'noexcept}}
 #if __cplusplus <= 201402L

clang/test/CXX/drs/dr12xx.cpp

@@ -3,7 +3,7 @@
 // RUN: %clang_cc1 -std=c++14 %s -verify -fexceptions -fcxx-exceptions -pedantic-errors
 // RUN: %clang_cc1 -std=c++1z %s -verify -fexceptions -fcxx-exceptions -pedantic-errors
 
-namespace dr1213 { // dr1213: 4.0
+namespace dr1213 { // dr1213: 4
 #if __cplusplus >= 201103L
   using T = int[3];
   int &&r = T{}[1];
@@ -26,7 +26,7 @@ struct Derived : Base {
 };
 } // dr1250
 
-namespace dr1295 {  // dr1295: 4.0
+namespace dr1295 {  // dr1295: 4
   struct X {
     unsigned bitfield : 4;
   };

clang/test/CXX/drs/dr13xx.cpp

@@ -31,7 +31,7 @@ namespace dr1315 { // dr1315: partial
   // expected-error@-1 {{type of specialized non-type template argument depends on a template parameter of the partial specialization}}
 }
 
-namespace dr1330 { // dr1330: 4.0 c++11
+namespace dr1330 { // dr1330: 4 c++11
   // exception-specifications are parsed in a context where the class is complete.
   struct A {
     void f() throw(T) {} // expected-error 0-1{{C++1z}} expected-note 0-1{{noexcept}}
@@ -175,7 +175,7 @@ namespace dr1359 { // dr1359: 3.5
 #endif
 }
 
-namespace dr1388 { // dr1388: 4.0
+namespace dr1388 { // dr1388: 4
   template<typename A, typename ...T> void f(T..., A); // expected-note 1+{{candidate}} expected-error 0-1{{C++11}}
   template<typename ...T> void g(T..., int); // expected-note 1+{{candidate}} expected-error 0-1{{C++11}}
   template<typename ...T, typename A> void h(T..., A); // expected-note 1+{{candidate}} expected-error 0-1{{C++11}}

clang/test/CXX/drs/dr14xx.cpp

@@ -343,7 +343,7 @@ namespace dr1490 {  // dr1490: 3.7 c++11
   std::initializer_list<char>{"abc"}; // expected-error {{expected unqualified-id}}}
 } // dr190
 
-namespace dr1495 { // dr1495: 4.0
+namespace dr1495 { // dr1495: 4
   // Deduction succeeds in both directions.
   template<typename T, typename U> struct A {}; // expected-note {{template is declared here}}
   template<typename T, typename U> struct A<U, T> {}; // expected-error {{class template partial specialization is not more specialized}}

clang/test/CXX/drs/dr15xx.cpp

@@ -3,7 +3,7 @@
 // RUN: %clang_cc1 -std=c++14 -triple x86_64-unknown-unknown %s -verify -fexceptions -fcxx-exceptions -pedantic-errors
 // RUN: %clang_cc1 -std=c++1z -triple x86_64-unknown-unknown %s -verify -fexceptions -fcxx-exceptions -pedantic-errors
 
-namespace dr1512 { // dr1512: 4.0
+namespace dr1512 { // dr1512: 4
   void f(char *p) {
     if (p > 0) {} // expected-error {{ordered comparison between pointer and zero}}
 #if __cplusplus >= 201103L
@@ -135,7 +135,7 @@ namespace dr1512 { // dr1512: 4.0
   }
 }
 
-namespace dr1518 { // dr1518: 4.0
+namespace dr1518 { // dr1518: 4
 #if __cplusplus >= 201103L
 struct Z0 { // expected-note 0+ {{candidate}}
   explicit Z0() = default; // expected-note 0+ {{here}}

clang/test/CXX/drs/dr16xx.cpp

@@ -86,7 +86,7 @@ namespace dr1645 { // dr1645: 3.9
 #endif
 }
 
-namespace dr1653 { // dr1653: 4.0 c++17
+namespace dr1653 { // dr1653: 4 c++17
   void f(bool b) {
     ++b;
     b++;

clang/test/CXX/drs/dr18xx.cpp

@@ -7,7 +7,7 @@
 // expected-no-diagnostics
 #endif
 
-void dr1891() { // dr1891: 4.0
+void dr1891() { // dr1891: 4
 #if __cplusplus >= 201103L
   int n;
   auto a = []{}; // expected-note 2{{candidate}} expected-note 2{{here}}

clang/test/CXX/drs/dr2xx.cpp

@@ -679,7 +679,7 @@ namespace dr258 { // dr258: yes
   } f; // expected-error {{abstract}}
 }
 
-namespace dr259 { // dr259: 4.0
+namespace dr259 { // dr259: 4
   template<typename T> struct A {};
   template struct A<int>; // expected-note {{previous}}
   template struct A<int>; // expected-error {{duplicate explicit instantiation}}

clang/test/CXX/drs/dr5xx.cpp

@@ -863,7 +863,7 @@ namespace dr580 { // dr580: partial
 
 // dr582: na
 
-namespace dr583 { // dr583: 4.0
+namespace dr583 { // dr583: 4
   // see n3624
   int *p;
   bool b1 = p < 0; // expected-error {{ordered comparison between pointer and zero}}

clang/test/CXX/drs/dr6xx.cpp

@@ -142,7 +142,7 @@ namespace dr615 { // dr615: yes
   static int n = f();
 }
 
-namespace dr616 { // dr616: 4.0
+namespace dr616 { // dr616: 4
 #if __cplusplus >= 201103L
   struct S { int n; } s;
   S f();

clang/test/CodeCompletion/member-access.cpp

@@ -27,6 +27,16 @@ public:
 
 void test(const Proxy &p) {
   p->
+}
+
+struct Test1 {
+  Base1 b;
+
+  static void sfunc() {
+    b. // expected-error {{invalid use of member 'b' in static member function}}
+  }
+};
+
   // RUN: %clang_cc1 -fsyntax-only -code-completion-at=%s:29:6 %s -o - | FileCheck -check-prefix=CHECK-CC1 %s
   // CHECK-CC1: Base1 : Base1::
   // CHECK-CC1: member1 : [#int#][#Base1::#]member1
@@ -39,4 +49,6 @@ void test(const Proxy &p) {
   // CHECK-CC1: memfun1 (Hidden) : [#void#]Base2::memfun1(<#int#>)
   // CHECK-CC1: memfun2 : [#void#][#Base3::#]memfun2(<#int#>)
   // CHECK-CC1: memfun3 : [#int#]memfun3(<#int#>)
-  
+
+// Make sure this doesn't crash
+// RUN: %clang_cc1 -fsyntax-only -code-completion-at=%s:36:7 %s -verify

clang/test/CodeGen/always_inline.c

@@ -1,7 +1,7 @@
 // RUN: %clang -emit-llvm -S -o %t %s
 // RUN: not grep '@f0' %t
 // RUN: not grep 'call ' %t
-// RUN: %clang -mllvm -disable-llvm-passes -emit-llvm -S -o %t %s
+// RUN: %clang -Xclang -disable-llvm-passes -emit-llvm -S -o %t %s
 // RUN: grep '@f0' %t | count 2
 
 //static int f0() { 

clang/test/CodeGenCXX/cxx11-thread-local.cpp

@@ -6,18 +6,18 @@
 int f();
 int g();
 
-// LINUX: @a = thread_local global i32 0
+// LINUX-DAG: @a = thread_local global i32 0
-// DARWIN: @a = internal thread_local global i32 0
+// DARWIN-DAG: @a = internal thread_local global i32 0
 thread_local int a = f();
 extern thread_local int b;
-// CHECK: @c = global i32 0
+// CHECK-DAG: @c = global i32 0
 int c = b;
-// CHECK: @_ZL1d = internal thread_local global i32 0
+// CHECK-DAG: @_ZL1d = internal thread_local global i32 0
 static thread_local int d = g();
 
 struct U { static thread_local int m; };
-// LINUX: @_ZN1U1mE = thread_local global i32 0
+// LINUX-DAG: @_ZN1U1mE = thread_local global i32 0
-// DARWIN: @_ZN1U1mE = internal thread_local global i32 0
+// DARWIN-DAG: @_ZN1U1mE = internal thread_local global i32 0
 thread_local int U::m = f();
 
 namespace MismatchedInitType {
@@ -35,37 +35,64 @@ namespace MismatchedInitType {
 template<typename T> struct V { static thread_local int m; };
 template<typename T> thread_local int V<T>::m = g();
 
-// CHECK: @e = global i32 0
+template<typename T> struct W { static thread_local int m; };
-int e = V<int>::m;
+template<typename T> thread_local int W<T>::m = 123;
 
-// CHECK: @_ZN1VIiE1mE = linkonce_odr thread_local global i32 0
+struct Dtor { ~Dtor(); };
+template<typename T> struct X { static thread_local Dtor m; };
+template<typename T> thread_local Dtor X<T>::m;
 
-// CHECK: @_ZZ1fvE1n = internal thread_local global i32 0
+// CHECK-DAG: @e = global
+void *e = V<int>::m + W<int>::m + &X<int>::m;
 
-// CHECK: @_ZGVZ1fvE1n = internal thread_local global i8 0
+template thread_local int V<float>::m;
+template thread_local int W<float>::m;
+template thread_local Dtor X<float>::m;
 
-// CHECK: @_ZZ8tls_dtorvE1s = internal thread_local global
+extern template thread_local int V<char>::m;
-// CHECK: @_ZGVZ8tls_dtorvE1s = internal thread_local global i8 0
+extern template thread_local int W<char>::m;
+extern template thread_local Dtor X<char>::m;
 
-// CHECK: @_ZZ8tls_dtorvE1t = internal thread_local global
+void *e2 = V<char>::m + W<char>::m + &X<char>::m;
-// CHECK: @_ZGVZ8tls_dtorvE1t = internal thread_local global i8 0
 
-// CHECK: @_ZZ8tls_dtorvE1u = internal thread_local global
+// CHECK-DAG: @_ZN1VIiE1mE = linkonce_odr thread_local global i32 0
-// CHECK: @_ZGVZ8tls_dtorvE1u = internal thread_local global i8 0
+// CHECK-DAG: @_ZN1WIiE1mE = linkonce_odr thread_local global i32 123
-// CHECK: @_ZGRZ8tls_dtorvE1u_ = internal thread_local global
+// CHECK-DAG: @_ZN1XIiE1mE = linkonce_odr thread_local global {{.*}}
+// CHECK-DAG: @_ZN1VIfE1mE = weak_odr thread_local global i32 0
+// CHECK-DAG: @_ZN1WIfE1mE = weak_odr thread_local global i32 123
+// CHECK-DAG: @_ZN1XIfE1mE = weak_odr thread_local global {{.*}}
 
-// CHECK: @_ZGVN1VIiE1mE = linkonce_odr thread_local global i64 0
+// CHECK-DAG: @_ZZ1fvE1n = internal thread_local global i32 0
 
-// CHECK: @__tls_guard = internal thread_local global i8 0
+// CHECK-DAG: @_ZGVZ1fvE1n = internal thread_local global i8 0
 
-// CHECK: @llvm.global_ctors = appending global {{.*}} @[[GLOBAL_INIT:[^ ]*]]
+// CHECK-DAG: @_ZZ8tls_dtorvE1s = internal thread_local global
+// CHECK-DAG: @_ZGVZ8tls_dtorvE1s = internal thread_local global i8 0
 
-// LINUX: @_ZTH1a = alias void (), void ()* @__tls_init
+// CHECK-DAG: @_ZZ8tls_dtorvE1t = internal thread_local global
-// DARWIN: @_ZTH1a = internal alias void (), void ()* @__tls_init
+// CHECK-DAG: @_ZGVZ8tls_dtorvE1t = internal thread_local global i8 0
-// CHECK: @_ZTHL1d = internal alias void (), void ()* @__tls_init
+
-// LINUX: @_ZTHN1U1mE = alias void (), void ()* @__tls_init
+// CHECK-DAG: @_ZZ8tls_dtorvE1u = internal thread_local global
-// DARWIN: @_ZTHN1U1mE = internal alias void (), void ()* @__tls_init
+// CHECK-DAG: @_ZGVZ8tls_dtorvE1u = internal thread_local global i8 0
-// CHECK: @_ZTHN1VIiE1mE = linkonce_odr alias void (), void ()* @__tls_init
+// CHECK-DAG: @_ZGRZ8tls_dtorvE1u_ = internal thread_local global
+
+// CHECK-DAG: @_ZGVN1VIiE1mE = linkonce_odr thread_local global i64 0
+
+// CHECK-DAG: @__tls_guard = internal thread_local global i8 0
+
+// CHECK-DAG: @llvm.global_ctors = appending global {{.*}} @[[GLOBAL_INIT:[^ ]*]]
+
+// LINUX-DAG: @_ZTH1a = alias void (), void ()* @__tls_init
+// DARWIN-DAG: @_ZTH1a = internal alias void (), void ()* @__tls_init
+// CHECK-DAG: @_ZTHL1d = internal alias void (), void ()* @__tls_init
+// LINUX-DAG: @_ZTHN1U1mE = alias void (), void ()* @__tls_init
+// DARWIN-DAG: @_ZTHN1U1mE = internal alias void (), void ()* @__tls_init
+// CHECK-DAG: @_ZTHN1VIiE1mE = linkonce_odr alias void (), void ()* @[[V_M_INIT:[^, ]*]]
+// CHECK-NOT: @_ZTHN1WIiE1mE =
+// CHECK-DAG: @_ZTHN1XIiE1mE = linkonce_odr alias void (), void ()* @[[X_M_INIT:[^, ]*]]
+// CHECK-DAG: @_ZTHN1VIfE1mE = weak_odr alias void (), void ()* @[[VF_M_INIT:[^, ]*]]
+// CHECK-NOT: @_ZTHN1WIfE1mE =
+// CHECK-DAG: @_ZTHN1XIfE1mE = weak_odr alias void (), void ()* @[[XF_M_INIT:[^, ]*]]
 
 
 // Individual variable initialization functions:
@@ -118,7 +145,9 @@ int f() {
 // LINUX: call i32* @_ZTWN1VIiE1mE()
 // DARWIN: call cxx_fast_tlscc i32* @_ZTWN1VIiE1mE()
 // CHECK-NEXT: load i32, i32* %{{.*}}, align 4
-// CHECK-NEXT: store i32 %{{.*}}, i32* @e, align 4
+// LINUX: call {{.*}}* @_ZTWN1XIiE1mE()
+// DARWIN: call cxx_fast_tlscc {{.*}}* @_ZTWN1XIiE1mE()
+// CHECK: store {{.*}} @e
 
 // LINUX-LABEL: define weak_odr hidden i32* @_ZTWN1VIiE1mE()
 // DARWIN-LABEL: define weak_odr hidden cxx_fast_tlscc i32* @_ZTWN1VIiE1mE()
@@ -126,6 +155,64 @@ int f() {
 // DARWIN: call cxx_fast_tlscc void @_ZTHN1VIiE1mE()
 // CHECK: ret i32* @_ZN1VIiE1mE
 
+// LINUX-LABEL: define weak_odr hidden i32* @_ZTWN1WIiE1mE()
+// DARWIN-LABEL: define weak_odr hidden cxx_fast_tlscc i32* @_ZTWN1WIiE1mE()
+// CHECK-NOT: call
+// CHECK: ret i32* @_ZN1WIiE1mE
+
+// LINUX-LABEL: define weak_odr hidden {{.*}}* @_ZTWN1XIiE1mE()
+// DARWIN-LABEL: define weak_odr hidden cxx_fast_tlscc {{.*}}* @_ZTWN1XIiE1mE()
+// LINUX: call void @_ZTHN1XIiE1mE()
+// DARWIN: call cxx_fast_tlscc void @_ZTHN1XIiE1mE()
+// CHECK: ret {{.*}}* @_ZN1XIiE1mE
+
+// CHECK: define internal {{.*}} @[[VF_M_INIT]]()
+// LINUX-SAME: comdat($_ZN1VIfE1mE)
+// DARWIN-NOT: comdat
+// CHECK: load i8, i8* bitcast (i64* @_ZGVN1VIfE1mE to i8*)
+// CHECK: %[[VF_M_INITIALIZED:.*]] = icmp eq i8 %{{.*}}, 0
+// CHECK: br i1 %[[VF_M_INITIALIZED]],
+// need init:
+// CHECK: call i32 @_Z1gv()
+// CHECK: store i32 %{{.*}}, i32* @_ZN1VIfE1mE, align 4
+// CHECK: store i64 1, i64* @_ZGVN1VIfE1mE
+// CHECK: br label
+
+// CHECK: define internal {{.*}} @[[XF_M_INIT]]()
+// LINUX-SAME: comdat($_ZN1XIfE1mE)
+// DARWIN-NOT: comdat
+// CHECK: load i8, i8* bitcast (i64* @_ZGVN1XIfE1mE to i8*)
+// CHECK: %[[XF_M_INITIALIZED:.*]] = icmp eq i8 %{{.*}}, 0
+// CHECK: br i1 %[[XF_M_INITIALIZED]],
+// need init:
+// LINUX: call {{.*}}__cxa_thread_atexit
+// DARWIN: call {{.*}}_tlv_atexit
+// CHECK: store i64 1, i64* @_ZGVN1XIfE1mE
+// CHECK: br label
+
+// LINUX: declare i32 @__cxa_thread_atexit(void (i8*)*, i8*, i8*)
+// DARWIN: declare i32 @_tlv_atexit(void (i8*)*, i8*, i8*)
+
+// DARWIN: declare cxx_fast_tlscc i32* @_ZTWN1VIcE1mE()
+// LINUX: define weak_odr hidden i32* @_ZTWN1VIcE1mE()
+// LINUX-NOT: comdat
+// LINUX: br i1 icmp ne (void ()* @_ZTHN1VIcE1mE,
+// LINUX: call void @_ZTHN1VIcE1mE()
+// LINUX: ret i32* @_ZN1VIcE1mE
+
+// DARWIN: declare cxx_fast_tlscc i32* @_ZTWN1WIcE1mE()
+// LINUX: define weak_odr hidden i32* @_ZTWN1WIcE1mE()
+// LINUX-NOT: comdat
+// LINUX: br i1 icmp ne (void ()* @_ZTHN1WIcE1mE,
+// LINUX: call void @_ZTHN1WIcE1mE()
+// LINUX: ret i32* @_ZN1WIcE1mE
+
+// DARWIN: declare cxx_fast_tlscc {{.*}}* @_ZTWN1XIcE1mE()
+// LINUX: define weak_odr hidden {{.*}}* @_ZTWN1XIcE1mE()
+// LINUX-NOT: comdat
+// LINUX: br i1 icmp ne (void ()* @_ZTHN1XIcE1mE,
+// LINUX: call void @_ZTHN1XIcE1mE()
+// LINUX: ret {{.*}}* @_ZN1XIcE1mE
 
 struct S { S(); ~S(); };
 struct T { ~T(); };
@@ -154,9 +241,6 @@ void tls_dtor() {
   static thread_local const S &u = S();
 }
 
-// LINUX: declare i32 @__cxa_thread_atexit(void (i8*)*, i8*, i8*)
-// DARWIN: declare i32 @_tlv_atexit(void (i8*)*, i8*, i8*)
-
 // CHECK: define {{.*}} @_Z7PR15991v(
 int PR15991() {
   thread_local int n;
@@ -184,7 +268,9 @@ void set_anon_i() {
 // LINUX-LABEL: define internal i32* @_ZTWN12_GLOBAL__N_16anon_iE()
 // DARWIN-LABEL: define internal cxx_fast_tlscc i32* @_ZTWN12_GLOBAL__N_16anon_iE()
 
-// CHECK: define {{.*}} @[[V_M_INIT:.*]]()
+// CHECK: define internal {{.*}} @[[V_M_INIT]]()
+// LINUX-SAME: comdat($_ZN1VIiE1mE)
+// DARWIN-NOT: comdat
 // CHECK: load i8, i8* bitcast (i64* @_ZGVN1VIiE1mE to i8*)
 // CHECK: %[[V_M_INITIALIZED:.*]] = icmp eq i8 %{{.*}}, 0
 // CHECK: br i1 %[[V_M_INITIALIZED]],
@@ -194,6 +280,18 @@ void set_anon_i() {
 // CHECK: store i64 1, i64* @_ZGVN1VIiE1mE
 // CHECK: br label
 
+// CHECK: define internal {{.*}} @[[X_M_INIT]]()
+// LINUX-SAME: comdat($_ZN1XIiE1mE)
+// DARWIN-NOT: comdat
+// CHECK: load i8, i8* bitcast (i64* @_ZGVN1XIiE1mE to i8*)
+// CHECK: %[[X_M_INITIALIZED:.*]] = icmp eq i8 %{{.*}}, 0
+// CHECK: br i1 %[[X_M_INITIALIZED]],
+// need init:
+// LINUX: call {{.*}}__cxa_thread_atexit
+// DARWIN: call {{.*}}_tlv_atexit
+// CHECK: store i64 1, i64* @_ZGVN1XIiE1mE
+// CHECK: br label
+
 // CHECK: define {{.*}}@[[GLOBAL_INIT:.*]]()
 // CHECK: call void @[[C_INIT]]()
 // CHECK: call void @[[E_INIT]]()
@@ -205,10 +303,13 @@ void set_anon_i() {
 // CHECK: br i1 %[[NEED_TLS_INIT]],
 // init:
 // CHECK: store i8 1, i8* @__tls_guard
+// CHECK-NOT: call void @[[V_M_INIT]]()
 // CHECK: call void @[[A_INIT]]()
+// CHECK-NOT: call void @[[V_M_INIT]]()
 // CHECK: call void @[[D_INIT]]()
+// CHECK-NOT: call void @[[V_M_INIT]]()
 // CHECK: call void @[[U_M_INIT]]()
-// CHECK: call void @[[V_M_INIT]]()
+// CHECK-NOT: call void @[[V_M_INIT]]()
 
 
 // LIUNX: define weak_odr hidden i32* @_ZTW1a() {

clang/test/CodeGenCXX/dllexport.cpp

@@ -732,13 +732,27 @@ USEMEMFUNC(ExplicitInstantiationDeclExportedDefTemplate<int>, f);
 // M32-DAG: define weak_odr dllexport x86_thiscallcc %struct.ExplicitInstantiationDeclExportedDefTemplate* @"\01??0?$ExplicitInstantiationDeclExportedDefTemplate@H@@QAE@XZ"
 // G32-DAG: define weak_odr x86_thiscallcc void @_ZN44ExplicitInstantiationDeclExportedDefTemplateIiE1fEv
 
-template <typename T> struct ImplicitInstantiationExplicitInstantiationDefExportedTemplate { void f() {} };
+template <typename T> struct ImplicitInstantiationExportedExplicitInstantiationDefTemplate { virtual void f() {} };
+ImplicitInstantiationExportedExplicitInstantiationDefTemplate<int> ImplicitInstantiationExportedExplicitInstantiationDefTemplateInstance;
+template struct __declspec(dllexport) ImplicitInstantiationExportedExplicitInstantiationDefTemplate<int>;
+USEMEMFUNC(ImplicitInstantiationExportedExplicitInstantiationDefTemplate<int>, f);
+// M32-DAG: define weak_odr dllexport x86_thiscallcc void @"\01?f@?$ImplicitInstantiationExportedExplicitInstantiationDefTemplate@H@@UAEXXZ"
+// G32-DAG: define weak_odr x86_thiscallcc void @_ZN61ImplicitInstantiationExportedExplicitInstantiationDefTemplateIiE1fEv
+
+template <typename T> struct __declspec(dllexport) ImplicitInstantiationExplicitInstantiationDefExportedTemplate { virtual void f() {} };
 ImplicitInstantiationExplicitInstantiationDefExportedTemplate<int> ImplicitInstantiationExplicitInstantiationDefExportedTemplateInstance;
-template class __declspec(dllexport) ImplicitInstantiationExplicitInstantiationDefExportedTemplate<int>;
+template struct ImplicitInstantiationExplicitInstantiationDefExportedTemplate<int>;
 USEMEMFUNC(ImplicitInstantiationExplicitInstantiationDefExportedTemplate<int>, f);
-// M32-DAG: define weak_odr dllexport x86_thiscallcc void @"\01?f@?$ImplicitInstantiationExplicitInstantiationDefExportedTemplate@H@@QAEXXZ"
+// M32-DAG: define weak_odr dllexport x86_thiscallcc void @"\01?f@?$ImplicitInstantiationExplicitInstantiationDefExportedTemplate@H@@UAEXXZ"
 // G32-DAG: define weak_odr x86_thiscallcc void @_ZN61ImplicitInstantiationExplicitInstantiationDefExportedTemplateIiE1fEv
 
+template <typename T> struct __declspec(dllexport) ImplicitInstantiationExportedExplicitInstantiationDefExportedTemplate { virtual void f() {} };
+ImplicitInstantiationExportedExplicitInstantiationDefExportedTemplate<int> ImplicitInstantiationExportedExplicitInstantiationDefExportedTemplateInstance;
+template struct __declspec(dllexport) ImplicitInstantiationExportedExplicitInstantiationDefExportedTemplate<int>;
+USEMEMFUNC(ImplicitInstantiationExportedExplicitInstantiationDefExportedTemplate<int>, f);
+// M32-DAG: define weak_odr dllexport x86_thiscallcc void @"\01?f@?$ImplicitInstantiationExportedExplicitInstantiationDefExportedTemplate@H@@UAEXXZ"
+// G32-DAG: define weak_odr x86_thiscallcc void @_ZN69ImplicitInstantiationExportedExplicitInstantiationDefExportedTemplateIiE1fEv
+
 namespace { struct InternalLinkageType {}; }
 struct __declspec(dllexport) PR23308 {
   void f(InternalLinkageType*);

clang/test/Driver/cl-options.c

@@ -535,7 +535,7 @@
 // RUN:     -fno-ms-compatibility \
 // RUN:     -fms-extensions \
 // RUN:     -fno-ms-extensions \
-// RUN:     -mllvm -disable-llvm-passes \
+// RUN:     -Xclang -disable-llvm-passes \
 // RUN:     -resource-dir asdf \
 // RUN:     -resource-dir=asdf \
 // RUN:     -Wunused-variable \

clang/test/Driver/disable-llvm.c

@@ -0,0 +1,22 @@
+// We support a CC1 option for disabling LLVM's passes.
+// RUN: %clang -O2 -Xclang -disable-llvm-passes -### %s 2>&1 \
+// RUN:     | FileCheck --check-prefix=DISABLED %s
+// DISABLED: -cc1
+// DISABLED-NOT: "-mllvm" "-disable-llvm-passes"
+// DISABLED: "-disable-llvm-passes"
+//
+// We also support two alternative spellings for historical reasons.
+// RUN: %clang -O2 -Xclang -disable-llvm-optzns -### %s 2>&1 \
+// RUN:     | FileCheck --check-prefix=DISABLED-LEGACY %s
+// RUN: %clang -O2 -mllvm -disable-llvm-optzns -### %s 2>&1 \
+// RUN:     | FileCheck --check-prefix=DISABLED-LEGACY %s
+// DISABLED-LEGACY: -cc1
+// DISABLED-LEGACY-NOT: "-mllvm" "-disable-llvm-optzns"
+// DISABLED-LEGACY: "-disable-llvm-optzns"
+//
+// The main flag shouldn't be specially handled when used with '-mllvm'.
+// RUN: %clang -O2 -mllvm -disable-llvm-passes -### %s 2>&1 | FileCheck --check-prefix=MLLVM %s
+// MLLVM: -cc1
+// MLLVM-NOT: -disable-llvm-passes
+// MLLVM: "-mllvm" "-disable-llvm-passes"
+// MLLVM-NOT: -disable-llvm-passes

clang/test/OpenMP/threadprivate_codegen.cpp

@@ -179,9 +179,9 @@ struct S5 {
 // CHECK-TLS-DAG:  @__dso_handle = external global i8
 // CHECK-TLS-DAG:  [[GS1_TLS_INIT:@_ZTHL3gs1]] = internal alias void (), void ()* @__tls_init
 // CHECK-TLS-DAG:  [[ARR_X_TLS_INIT:@_ZTH5arr_x]] = alias void (), void ()* @__tls_init
-// CHECK-TLS-DAG:  [[ST_INT_ST_TLS_INIT:@_ZTHN2STIiE2stE]] = linkonce_odr alias void (), void ()* @__tls_init
+
-// CHECK-TLS-DAG:  [[ST_FLOAT_ST_TLS_INIT:@_ZTHN2STIfE2stE]] = linkonce_odr alias void (), void ()* @__tls_init
+
-// CHECK-TLS-DAG:  [[ST_S4_ST_TLS_INIT:@_ZTHN2STI2S4E2stE]] = linkonce_odr alias void (), void ()* @__tls_init
+// CHECK-TLS-DAG:  [[ST_S4_ST_TLS_INIT:@_ZTHN2STI2S4E2stE]] = linkonce_odr alias void (), void ()* [[ST_S4_ST_CXX_INIT:@[^, ]*]]
 
 struct Static {
   static S3 s;
@@ -640,11 +640,11 @@ int main() {
 // CHECK-TLS:   ret [2 x [3 x [[S1]]]]* [[ARR_X]]
 // CHECK-TLS: }
 // CHECK-TLS: define {{.*}} i32* [[ST_INT_ST_TLS_INITD]] {{#[0-9]+}} {
-// CHECK-TLS:   call void [[ST_INT_ST_TLS_INIT]]
+// CHECK-TLS-NOT:   call
 // CHECK-TLS:   ret i32* [[ST_INT_ST]]
 // CHECK-TLS: }
 // CHECK-TLS: define {{.*}} float* [[ST_FLOAT_ST_TLS_INITD]] {{#[0-9]+}} {
-// CHECK-TLS:   call void [[ST_FLOAT_ST_TLS_INIT]]
+// CHECK-TLS-NOT:   call
 // CHECK-TLS:   ret float* [[ST_FLOAT_ST]]
 // CHECK-TLS: }
 // CHECK-TLS: define {{.*}} [[S4]]* [[ST_S4_ST_TLS_INITD]] {{#[0-9]+}} {
@@ -923,7 +923,7 @@ int foobar() {
 // CHECK-TLS: define {{.*}}void [[SM_CTOR2]]([[SMAIN]]* {{.*}}, i32 {{.*}})
 // CHECK-TLS: define {{.*}}void [[SM_DTOR2]]([[SMAIN]]* {{.*}})
 
-// CHECK-TLS: define internal void [[ST_S4_ST_CXX_INIT:@.*]]()
+// CHECK-TLS: define internal void [[ST_S4_ST_CXX_INIT]]()
 // CHECK-TLS: call void [[ST_S4_ST_CTOR1:@.*]]([[S4]]* [[ST_S4_ST]], i32 23)
 // CHECK-TLS: call i32 @__cxa_thread_atexit(void (i8*)* bitcast (void ([[S4]]*)* [[ST_S4_ST_DTOR1:.*]] to void (i8*)*), i8* bitcast ([[S4]]* [[ST_S4_ST]] to i8*)
 // CHECK-TLS: }
@@ -945,7 +945,7 @@ int foobar() {
 // CHECK-TLS:      call void [[GS1_CXX_INIT]]
 // CHECK-TLS-NOT:  call void [[GS2_CXX_INIT]]
 // CHECK-TLS:      call void [[ARR_X_CXX_INIT]]
-// CHECK-TLS:      call void [[ST_S4_ST_CXX_INIT]]
+// CHECK-TLS-NOT:  call void [[ST_S4_ST_CXX_INIT]]
 // CHECK-TLS:      [[DONE_LABEL]]
 
 // CHECK-TLS-DAG:      declare {{.*}} void [[GS3_TLS_INIT]]

clang/test/Preprocessor/dependencies-and-pp.c

@@ -32,5 +32,12 @@
 // RUN: FileCheck -check-prefix=TEST5 %s < %t.d
 // TEST5: foo $$(bar) b az qu\ ux \ space:
 
+// Test self dependency, PR31644
+
+// RUN: %clang -E -MD -MP -MF %t.d %s
+// RUN: FileCheck -check-prefix=TEST6 %s < %t.d
+// TEST6: dependencies-and-pp.c
+// TEST6-NOT: dependencies-and-pp.c:
+
 // TODO: Test default target without quoting
 // TODO: Test default target with quoting

clang/test/Profile/gcc-flag-compatibility.c

@@ -18,14 +18,14 @@
 // RUN: rm -rf %t.dir
 // RUN: mkdir -p %t.dir/some/path
 // RUN: llvm-profdata merge %S/Inputs/gcc-flag-compatibility.proftext -o %t.dir/some/path/default.profdata
-// RUN: %clang %s -o - -mllvm -disable-llvm-passes -emit-llvm -S -fprofile-use=%t.dir/some/path | FileCheck -check-prefix=PROFILE-USE-2 %s
+// RUN: %clang %s -o - -Xclang -disable-llvm-passes -emit-llvm -S -fprofile-use=%t.dir/some/path | FileCheck -check-prefix=PROFILE-USE-2 %s
 // PROFILE-USE-2: = !{!"branch_weights", i32 101, i32 2}
 
 // Check that -fprofile-use=some/path/file.prof reads some/path/file.prof
 // RUN: rm -rf %t.dir
 // RUN: mkdir -p %t.dir/some/path
 // RUN: llvm-profdata merge %S/Inputs/gcc-flag-compatibility.proftext -o %t.dir/some/path/file.prof
-// RUN: %clang %s -o - -mllvm -disable-llvm-passes -emit-llvm -S -fprofile-use=%t.dir/some/path/file.prof | FileCheck -check-prefix=PROFILE-USE-3 %s
+// RUN: %clang %s -o - -Xclang -disable-llvm-passes -emit-llvm -S -fprofile-use=%t.dir/some/path/file.prof | FileCheck -check-prefix=PROFILE-USE-3 %s
 // PROFILE-USE-3: = !{!"branch_weights", i32 101, i32 2}
 
 int X = 0;

clang/test/SemaCXX/uninitialized.cpp

@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsyntax-only -Wall -Wuninitialized -Wno-unused-value -std=c++11 -verify %s
+// RUN: %clang_cc1 -fsyntax-only -Wall -Wuninitialized -Wno-unused-value -std=c++1z -verify %s
 
 // definitions for std::move
 namespace std {
@@ -1437,3 +1437,13 @@ void array_capture(bool b) {
     [fname]{};
   }
 }
+
+void if_switch_init_stmt(int k) {
+  if (int n = 0; (n == k || k > 5)) {}
+
+  if (int n; (n == k || k > 5)) {} // expected-warning {{uninitialized}} expected-note {{initialize}}
+
+  switch (int n = 0; (n == k || k > 5)) {} // expected-warning {{boolean}}
+
+  switch (int n; (n == k || k > 5)) {} // expected-warning {{uninitialized}} expected-note {{initialize}} expected-warning {{boolean}}
+}

clang/test/SemaTemplate/class-template-spec.cpp

@@ -207,19 +207,19 @@ namespace NTTPTypeVsPartialOrder {
   struct X { typedef int value_type; };
   template<typename T> struct Y { typedef T value_type; };
 
-  template<typename T, typename T::value_type N> struct A; // expected-note {{template}}
+  template<typename T, typename T::value_type N> struct A;
   template<int N> struct A<X, N> {};
-  template<typename T, T N> struct A<Y<T>, N> {}; // expected-error {{not more specialized}} expected-note {{'T' vs 'typename Y<type-parameter-0-0>::value_type'}}
+  template<typename T, T N> struct A<Y<T>, N> {};
   A<X, 0> ax;
   A<Y<int>, 0> ay;
 
 
-  template<int, typename T, typename T::value_type> struct B; // expected-note {{template}}
+  template<int, typename T, typename T::value_type> struct B;
-  template<typename T, typename T::value_type N> struct B<0, T, N>; // expected-note {{matches}}
+  template<typename T, typename T::value_type N> struct B<0, T, N>;
   template<int N> struct B<0, X, N> {};
-  template<typename T, T N> struct B<0, Y<T>, N> {}; // expected-error {{not more specialized}} expected-note {{'T' vs 'typename Y<type-parameter-0-0>::value_type'}} expected-note {{matches}}
+  template<typename T, T N> struct B<0, Y<T>, N> {};
   B<0, X, 0> bx;
-  B<0, Y<int>, 0> by; // expected-error {{ambiguous}}
+  B<0, Y<int>, 0> by;
 }
 
 namespace DefaultArgVsPartialSpec {

clang/test/SemaTemplate/partial-order.cpp

@@ -0,0 +1,14 @@
+// RUN: %clang_cc1 -std=c++1z %s -verify
+
+// expected-no-diagnostics
+
+namespace hana_enable_if_idiom {
+  template<bool> struct A {};
+  template<typename, typename = A<true>> struct B;
+  template<typename T, bool N> struct B<T, A<N>> {};
+  template<typename T> struct B<T, A<T::value>> {};
+  struct C {
+    static const bool value = true;
+  };
+  B<C> b;
+}

clang/test/SemaTemplate/temp_arg_nontype.cpp

@@ -370,13 +370,13 @@ namespace PR17696 {
 }
 
 namespace partial_order_different_types {
-  // These are unordered because the type of the final argument doesn't match.
+  template<int, int, typename T, typename, T> struct A;
-  template<int, int, typename T, typename, T> struct A; // expected-note {{here}}
+  template<int N, typename T, typename U, T V> struct A<0, N, T, U, V>; // expected-note {{matches}}
-  template<int N, typename T, typename U, T V> struct A<0, N, T, U, V> {}; // expected-note {{matches}}
+  // FIXME: It appears that this partial specialization should be ill-formed as
-  template<typename T, typename U, U V> struct A<0, 0, T, U, V> {}; // expected-note {{matches}}
+  // it is not more specialized than the primary template. V is not deducible
-  // expected-error@-1 {{not more specialized than the primary}}
+  // because it does not have the same type as the corresponding parameter.
-  // expected-note@-2 {{deduced non-type template argument does not have the same type as the corresponding template parameter ('U' vs 'type-parameter-0-0')}}
+  template<int N, typename T, typename U, U V> struct A<0, N, T, U, V> {}; // expected-note {{matches}}
-  A<0, 0, int, int, 0> a; // expected-error {{ambiguous partial specializations}}
+  A<0, 0, int, int, 0> a; // expected-error {{ambiguous}}
 }
 
 namespace partial_order_references {
@@ -434,7 +434,7 @@ namespace dependent_nested_partial_specialization {
 
   template<typename T> struct E {
     template<typename U, U V> struct F; // expected-note {{template}}
-    template<typename W, T V> struct F<W, V> {}; // expected-error {{not more specialized than the primary}} expected-note {{does not have the same type}}
+    template<typename W, T V> struct F<W, V> {}; // expected-error {{not more specialized than the primary}}
   };
   E<int>::F<int, 0> e1; // expected-note {{instantiation of}}
 }

clang/test/SemaTemplate/temp_arg_template_cxx1z.cpp

@@ -79,13 +79,13 @@ namespace Auto {
 
   TInt<Auto> ia;
   TInt<AutoPtr> iap; // expected-error {{different template parameters}}
-  TInt<DecltypeAuto> ida; // FIXME expected-error {{different template parameters}}
+  TInt<DecltypeAuto> ida;
   TInt<Int> ii;
   TInt<IntPtr> iip; // expected-error {{different template parameters}}
 
   TIntPtr<Auto> ipa;
   TIntPtr<AutoPtr> ipap;
-  TIntPtr<DecltypeAuto> ipda; // FIXME expected-error {{different template parameters}}
+  TIntPtr<DecltypeAuto> ipda;
   TIntPtr<Int> ipi; // expected-error {{different template parameters}}
   TIntPtr<IntPtr> ipip;
 
@@ -114,6 +114,6 @@ namespace Auto {
 
   int n;
   template<auto A, decltype(A) B = &n> struct SubstFailure;
-  TInt<SubstFailure> isf; // expected-error {{different template parameters}}
+  TInt<SubstFailure> isf; // FIXME: this should be ill-formed
-  TIntPtr<SubstFailure> ipsf; // expected-error {{different template parameters}}
+  TIntPtr<SubstFailure> ipsf;
 }

clang/www/cxx_dr_status.html

@@ -591,7 +591,7 @@
     <td><a href="http://www.open-std.org/jtc1/sc22/wg21/docs/cwg_defects.html#92">92</a></td>
     <td>WP</td>
     <td>Should <I>exception-specification</I>s be part of the type system?</td>
-    <td class="svn" align="center">SVN (C++17 onwards)</td>
+    <td class="svn" align="center">Clang 4 (C++17 onwards)</td>
   </tr>
   <tr id="93">
     <td><a href="http://www.open-std.org/jtc1/sc22/wg21/docs/cwg_defects.html#93">93</a></td>
@@ -1594,7 +1594,7 @@ accessible?</td>
     <td><a href="http://www.open-std.org/jtc1/sc22/wg21/docs/cwg_defects.html#259">259</a></td>
     <td>CD1</td>
     <td>Restrictions on explicit specialization and instantiation</td>
-    <td class="svn" align="center">SVN</td>
+    <td class="svn" align="center">Clang 4</td>
   </tr>
   <tr class="open" id="260">
     <td><a href="http://www.open-std.org/jtc1/sc22/wg21/docs/cwg_active.html#260">260</a></td>
@@ -3541,7 +3541,7 @@ and <I>POD class</I></td>
     <td><a href="http://www.open-std.org/jtc1/sc22/wg21/docs/cwg_defects.html#583">583</a></td>
     <td>CD3</td>
     <td>Relational pointer comparisons against the null pointer constant</td>
-    <td class="svn" align="center">SVN</td>
+    <td class="svn" align="center">Clang 4</td>
   </tr>
   <tr id="584">
     <td><a href="http://www.open-std.org/jtc1/sc22/wg21/docs/cwg_closed.html#584">584</a></td>
@@ -3739,7 +3739,7 @@ and <I>POD class</I></td>
     <td><a href="http://www.open-std.org/jtc1/sc22/wg21/docs/cwg_defects.html#616">616</a></td>
     <td>CD3</td>
     <td>Definition of &#8220;indeterminate value&#8221;</td>
-    <td class="svn" align="center">SVN</td>
+    <td class="svn" align="center">Clang 4</td>
   </tr>
   <tr class="open" id="617">
     <td><a href="http://www.open-std.org/jtc1/sc22/wg21/docs/cwg_active.html#617">617</a></td>
@@ -7093,7 +7093,7 @@ and <I>POD class</I></td>
     <td><a href="http://www.open-std.org/jtc1/sc22/wg21/docs/cwg_defects.html#1213">1213</a></td>
     <td>CD3</td>
     <td>Array subscripting and xvalues</td>
-    <td class="svn" align="center">SVN</td>
+    <td class="svn" align="center">Clang 4</td>
   </tr>
   <tr id="1214">
     <td><a href="http://www.open-std.org/jtc1/sc22/wg21/docs/cwg_defects.html#1214">1214</a></td>
@@ -7585,7 +7585,7 @@ and <I>POD class</I></td>
     <td><a href="http://www.open-std.org/jtc1/sc22/wg21/docs/cwg_defects.html#1295">1295</a></td>
     <td>CD3</td>
     <td>Binding a reference to an rvalue bit-field</td>
-    <td class="svn" align="center">SVN</td>
+    <td class="svn" align="center">Clang 4</td>
   </tr>
   <tr id="1296">
     <td><a href="http://www.open-std.org/jtc1/sc22/wg21/docs/cwg_defects.html#1296">1296</a></td>
@@ -7795,7 +7795,7 @@ and <I>POD class</I></td>
     <td><a href="http://www.open-std.org/jtc1/sc22/wg21/docs/cwg_defects.html#1330">1330</a></td>
     <td>CD3</td>
     <td>Delayed instantiation of <TT>noexcept</TT> specifiers</td>
-    <td class="svn" align="center">SVN (C++11 onwards)</td>
+    <td class="svn" align="center">Clang 4 (C++11 onwards)</td>
   </tr>
   <tr class="open" id="1331">
     <td><a href="http://www.open-std.org/jtc1/sc22/wg21/docs/cwg_closed.html#1331">1331</a></td>
@@ -8143,7 +8143,7 @@ and <I>POD class</I></td>
     <td><a href="http://www.open-std.org/jtc1/sc22/wg21/docs/cwg_defects.html#1388">1388</a></td>
     <td>CD3</td>
     <td>Missing non-deduced context following a function parameter pack</td>
-    <td class="svn" align="center">SVN</td>
+    <td class="svn" align="center">Clang 4</td>
   </tr>
   <tr id="1389">
     <td><a href="http://www.open-std.org/jtc1/sc22/wg21/docs/cwg_closed.html#1389">1389</a></td>
@@ -8785,7 +8785,7 @@ and <I>POD class</I></td>
     <td><a href="http://www.open-std.org/jtc1/sc22/wg21/docs/cwg_defects.html#1495">1495</a></td>
     <td>CD3</td>
     <td>Partial specialization of variadic class template</td>
-    <td class="svn" align="center">SVN</td>
+    <td class="svn" align="center">Clang 4</td>
   </tr>
   <tr id="1496">
     <td><a href="http://www.open-std.org/jtc1/sc22/wg21/docs/cwg_defects.html#1496">1496</a></td>
@@ -8887,7 +8887,7 @@ and <I>POD class</I></td>
     <td><a href="http://www.open-std.org/jtc1/sc22/wg21/docs/cwg_defects.html#1512">1512</a></td>
     <td>CD3</td>
     <td>Pointer comparison vs qualification conversions</td>
-    <td class="svn" align="center">SVN</td>
+    <td class="svn" align="center">Clang 4</td>
   </tr>
   <tr class="open" id="1513">
     <td><a href="http://www.open-std.org/jtc1/sc22/wg21/docs/cwg_active.html#1513">1513</a></td>
@@ -8923,7 +8923,7 @@ and <I>POD class</I></td>
     <td><a href="http://www.open-std.org/jtc1/sc22/wg21/docs/cwg_defects.html#1518">1518</a></td>
     <td>DRWP</td>
     <td>Explicit default constructors and copy-list-initialization</td>
-    <td class="svn" align="center">SVN</td>
+    <td class="svn" align="center">Clang 4</td>
   </tr>
   <tr id="1519">
     <td><a href="http://www.open-std.org/jtc1/sc22/wg21/docs/cwg_closed.html#1519">1519</a></td>
@@ -9733,7 +9733,7 @@ and <I>POD class</I></td>
     <td><a href="http://www.open-std.org/jtc1/sc22/wg21/docs/cwg_defects.html#1653">1653</a></td>
     <td>WP</td>
     <td>Removing deprecated increment of <TT>bool</TT></td>
-    <td class="svn" align="center">SVN (C++17 onwards)</td>
+    <td class="svn" align="center">Clang 4 (C++17 onwards)</td>
   </tr>
   <tr id="1654">
     <td><a href="http://www.open-std.org/jtc1/sc22/wg21/docs/cwg_closed.html#1654">1654</a></td>
@@ -11161,7 +11161,7 @@ and <I>POD class</I></td>
     <td><a href="http://www.open-std.org/jtc1/sc22/wg21/docs/cwg_defects.html#1891">1891</a></td>
     <td>DRWP</td>
     <td>Move constructor/assignment for closure class</td>
-    <td class="svn" align="center">SVN</td>
+    <td class="svn" align="center">Clang 4</td>
   </tr>
   <tr id="1892">
     <td><a href="http://www.open-std.org/jtc1/sc22/wg21/docs/cwg_defects.html#1892">1892</a></td>

clang/www/cxx_status.html

@@ -612,7 +612,7 @@ as the draft C++1z standard evolves.
     <tr>
       <td>Make exception specifications part of the type system</td>
       <td><a href="http://www.open-std.org/jtc1/sc22/wg21/docs/papers/2015/p0012r1.html">P0012R1</a></td>
-      <td class="svn" align="center">SVN</td>
+      <td class="svn" align="center">Clang 4</td>
     </tr>
     <tr>
       <td><tt>__has_include</tt> in preprocessor conditionals</td>
@@ -679,7 +679,7 @@ as the draft C++1z standard evolves.
     <tr>
       <td>Dynamic memory allocation for over-aligned data</td>
       <td><a href="http://wg21.link/p0035r4">P0035R4</a></td>
-      <td class="svn" align="center">SVN</td>
+      <td class="svn" align="center">Clang 4</td>
     </tr>
     <tr>
       <td>Template argument deduction for class templates</td>
@@ -689,17 +689,17 @@ as the draft C++1z standard evolves.
     <tr>
       <td>Non-type template parameters with <tt>auto</tt> type</td>
       <td><a href="http://wg21.link/p0127r2">P0127R2</a></td>
-      <td class="svn" align="center">SVN</td>
+      <td class="svn" align="center">Clang 4</td>
     </tr>
     <tr>
       <td>Guaranteed copy elision</td>
       <td><a href="http://wg21.link/p0135r1">P0135R1</a></td>
-      <td class="svn" align="center">SVN</td>
+      <td class="svn" align="center">Clang 4</td>
     </tr>
     <tr>
       <td rowspan=2>Stricter expression evaluation order</td>
       <td><a href="http://wg21.link/p0145r3">P0145R3</a></td>
-      <td class="svn" align="center" rowspan=2>SVN <a href="#p0145">(10)</a></td>
+      <td class="svn" align="center" rowspan=2>Clang 4 <a href="#p0145">(10)</a></td>
     </tr>
     <tr>
       <td><a href="http://wg21.link/p0400r0">P0400R0</a></td>
@@ -722,7 +722,7 @@ as the draft C++1z standard evolves.
     <tr>
       <td>Structured bindings</td>
       <td><a href="http://wg21.link/p0217r3">P0217R3</a></td>
-      <td class="svn" align="center">SVN</td>
+      <td class="svn" align="center">Clang 4</td>
     </tr>
     <tr>
       <td>Separate variable and condition for <tt>if</tt> and <tt>switch</tt></td>
@@ -738,12 +738,12 @@ as the draft C++1z standard evolves.
     <tr>
       <td>Removing deprecated dynamic exception specifications</td>
       <td><a href="http://wg21.link/p0003r5">P0003R5</a></td>
-      <td class="svn" align="center">SVN</td>
+      <td class="svn" align="center">Clang 4</td>
     </tr>
     <tr>
       <td>Pack expansions in <em>using-declarations</em></td>
       <td><a href="http://wg21.link/p0195r2">P0195R2</a></td>
-      <td class="svn" align="center">SVN</td>
+      <td class="svn" align="center">Clang 4</td>
     </tr>
 </table>
 
@@ -765,10 +765,10 @@ reverse construction order in that ABI.
 </span><br>
 <span id="p0522">(12): Despite being the the resolution to a Defect Report, this
 feature is disabled by default in all language versions, and can be enabled
-explicitly with the flag <tt>-frelaxed-template-template-args</tt>. The change
+explicitly with the flag <tt>-frelaxed-template-template-args</tt> in Clang 4.
-to the standard lacks a corresponding change for template partial ordering,
+The change to the standard lacks a corresponding change for template partial
-resulting in ambiguity errors for reasonable and previously-valid code. This
+ordering, resulting in ambiguity errors for reasonable and previously-valid
-issue is expected to be rectified soon.
+code. This issue is expected to be rectified soon.
 </span>
 </p>
 </details>
@@ -802,8 +802,8 @@ and library features that are not part of standard C++.</p>
       </td>
     </tr>
     <tr>
-      <td class="svn" align="center">
+      <td class="full" align="center">
-        SVN (<a href="http://wg21.link/p0096r3">P0096R3</a>)</a>
+        Clang 4 (<a href="http://wg21.link/p0096r3">P0096R3</a>)</a>
       </td>
     </tr>
     <!-- FIXME: Implement latest recommendations.

clang/www/make_cxx_dr_status

@@ -108,10 +108,13 @@ def availability(issue):
   if status == 'unknown':
     avail = 'Unknown'
     avail_style = ' class="none"'
-  elif status == '4.0':
+  elif status == '5':
     avail = 'SVN'
     avail_style = ' class="svn"'
-  elif re.match('^[0-9]+\.', status):
+  elif status == '4':
+    avail = 'Clang 4'
+    avail_style = ' class="svn"'
+  elif re.match('^[0-9]+\.?[0-9]*', status):
     avail = 'Clang %s' % status
     avail_style = ' class="full"'
   elif status == 'yes':

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/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,41 +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
-// RUN: %clang %target_itanium_abi_host_triple -arch x86_64 %t.o -o %t.out
-// RUN: %test_debuginfo %s %t.out
-//
-// 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,153 +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('^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/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/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;
-}

libclc/.gitignore

@@ -1,14 +0,0 @@
-Makefile
-amdgcn--
-amdgcn--amdhsa
-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