Feature/merge upstream 20210417

clang/docs/UsersManual.rst

@@ -686,7 +686,7 @@ generate a reproducer for warnings or errors while using modules.
 .. option:: -gen-reproducer
 
   Generates preprocessed source files, a reproducer script and if relevant, a
-  cache containing: built module pcm's and all headers needed to rebuilt the
+  cache containing: built module pcm's and all headers needed to rebuild the
   same modules.
 
 .. _rpass:

clang/include/clang/Basic/AttrDocs.td

@@ -3243,7 +3243,9 @@ def UnrollHintDocs : Documentation {
   let Content = [{
 Loop unrolling optimization hints can be specified with ``#pragma unroll`` and
 ``#pragma nounroll``. The pragma is placed immediately before a for, while,
-do-while, or c++11 range-based for loop.
+do-while, or c++11 range-based for loop. GCC's loop unrolling hints
+``#pragma GCC unroll`` and ``#pragma GCC nounroll`` are also supported and have
+identical semantics to ``#pragma unroll`` and ``#pragma nounroll``.
 
 Specifying ``#pragma unroll`` without a parameter directs the loop unroller to
 attempt to fully unroll the loop if the trip count is known at compile time and

clang/include/clang/Basic/TargetBuiltins.h

@@ -336,9 +336,9 @@ namespace clang {
   }
 
   static constexpr uint64_t LargestBuiltinID = std::max<uint64_t>(
-      {NEON::FirstTSBuiltin, ARM::LastTSBuiltin, SVE::FirstTSBuiltin,
-       AArch64::LastTSBuiltin, BPF::LastTSBuiltin, PPC::LastTSBuiltin,
-       NVPTX::LastTSBuiltin, AMDGPU::LastTSBuiltin, X86::LastTSBuiltin,
+      {ARM::LastTSBuiltin, AArch64::LastTSBuiltin, BPF::LastTSBuiltin,
+       PPC::LastTSBuiltin, NVPTX::LastTSBuiltin, AMDGPU::LastTSBuiltin,
+       X86::LastTSBuiltin, VE::LastTSBuiltin, RISCV::LastTSBuiltin,
        Hexagon::LastTSBuiltin, Mips::LastTSBuiltin, XCore::LastTSBuiltin,
        Le64::LastTSBuiltin, SystemZ::LastTSBuiltin,
        WebAssembly::LastTSBuiltin});

clang/include/clang/Tooling/NodeIntrospection.h

@@ -38,22 +38,16 @@ class LocationCall : public llvm::ThreadSafeRefCountedBase<LocationCall> {
   LocationCall(SharedLocationCall on, std::string name,
                LocationCallFlags flags = NoFlags)
       : m_flags(flags), m_on(std::move(on)), m_name(std::move(name)) {}
-  LocationCall(SharedLocationCall on, std::string name,
-               std::vector<std::string> args, LocationCallFlags flags = NoFlags)
-      : m_flags(flags), m_on(std::move(on)), m_name(std::move(name)),
-        m_args(std::move(args)) {}
 
   LocationCall *on() const { return m_on.get(); }
   StringRef name() const { return m_name; }
-  ArrayRef<std::string> args() const { return m_args; }
   bool returnsPointer() const { return m_flags & ReturnsPointer; }
   bool isCast() const { return m_flags & IsCast; }
 
 private:
   LocationCallFlags m_flags;
   SharedLocationCall m_on;
   std::string m_name;
-  std::vector<std::string> m_args;
 };
 
 class LocationCallFormatterCpp {
@@ -92,6 +86,7 @@ NodeLocationAccessors GetLocations(clang::CXXCtorInitializer const *Object);
 NodeLocationAccessors GetLocations(clang::NestedNameSpecifierLoc const *);
 NodeLocationAccessors GetLocations(clang::TemplateArgumentLoc const *);
 NodeLocationAccessors GetLocations(clang::CXXBaseSpecifier const *);
+NodeLocationAccessors GetLocations(clang::TypeLoc const &);
 NodeLocationAccessors GetLocations(clang::DynTypedNode const &Node);
 } // namespace NodeIntrospection
 } // namespace tooling

clang/lib/CodeGen/CGCUDANV.cpp

@@ -1089,6 +1089,28 @@ void CGNVCUDARuntime::transformManagedVars() {
 llvm::Function *CGNVCUDARuntime::finalizeModule() {
   if (CGM.getLangOpts().CUDAIsDevice) {
     transformManagedVars();
+
+    // Mark ODR-used device variables as compiler used to prevent it from being
+    // eliminated by optimization. This is necessary for device variables
+    // ODR-used by host functions. Sema correctly marks them as ODR-used no
+    // matter whether they are ODR-used by device or host functions.
+    //
+    // We do not need to do this if the variable has used attribute since it
+    // has already been added.
+    //
+    // Static device variables have been externalized at this point, therefore
+    // variables with LLVM private or internal linkage need not be added.
+    for (auto &&Info : DeviceVars) {
+      auto Kind = Info.Flags.getKind();
+      if (!Info.Var->isDeclaration() &&
+          !llvm::GlobalValue::isLocalLinkage(Info.Var->getLinkage()) &&
+          (Kind == DeviceVarFlags::Variable ||
+           Kind == DeviceVarFlags::Surface ||
+           Kind == DeviceVarFlags::Texture) &&
+          Info.D->isUsed() && !Info.D->hasAttr<UsedAttr>()) {
+        CGM.addCompilerUsedGlobal(Info.Var);
+      }
+    }
     return nullptr;
   }
   return makeModuleCtorFunction();

clang/lib/CodeGen/CGExprScalar.cpp

@@ -2099,6 +2099,7 @@ Value *ScalarExprEmitter::VisitCastExpr(CastExpr *CE) {
       DestLV.setTBAAInfo(TBAAAccessInfo::getMayAliasInfo());
       return EmitLoadOfLValue(DestLV, CE->getExprLoc());
     }
+
     return Builder.CreateBitCast(Src, DstTy);
   }
   case CK_AddressSpaceConversion: {

clang/lib/CodeGen/CodeGenFunction.cpp

@@ -174,7 +174,7 @@ void CodeGenFunction::CGFPOptionsRAII::ConstructorHelper(FPOptions FPFeatures) {
 
   auto mergeFnAttrValue = [&](StringRef Name, bool Value) {
     auto OldValue =
-        CGF.CurFn->getFnAttribute(Name).getValueAsString() == "true";
+        CGF.CurFn->getFnAttribute(Name).getValueAsBool();
     auto NewValue = OldValue & Value;
     if (OldValue != NewValue)
       CGF.CurFn->addFnAttr(Name, llvm::toStringRef(NewValue));

clang/lib/Headers/__clang_hip_cmath.h

@@ -14,13 +14,15 @@
 #error "This file is for HIP and OpenMP AMDGCN device compilation only."
 #endif
 
+#if !defined(__HIPCC_RTC__)
 #if defined(__cplusplus)
 #include <limits>
 #include <type_traits>
 #include <utility>
 #endif
 #include <limits.h>
 #include <stdint.h>
+#endif // __HIPCC_RTC__
 
 #pragma push_macro("__DEVICE__")
 #define __DEVICE__ static __device__ inline __attribute__((always_inline))

clang/lib/Headers/__clang_hip_math.h

@@ -13,11 +13,13 @@
 #error "This file is for HIP and OpenMP AMDGCN device compilation only."
 #endif
 
+#if !defined(__HIPCC_RTC__)
 #if defined(__cplusplus)
 #include <algorithm>
 #endif
 #include <limits.h>
 #include <stdint.h>
+#endif // __HIPCC_RTC__
 
 #pragma push_macro("__DEVICE__")
 #define __DEVICE__ static __device__ inline __attribute__((always_inline))
@@ -1260,13 +1262,15 @@ float min(float __x, float __y) { return fminf(__x, __y); }
 __DEVICE__
 double min(double __x, double __y) { return fmin(__x, __y); }
 
+#if !defined(__HIPCC_RTC__)
 __host__ inline static int min(int __arg1, int __arg2) {
   return std::min(__arg1, __arg2);
 }
 
 __host__ inline static int max(int __arg1, int __arg2) {
   return std::max(__arg1, __arg2);
 }
+#endif // __HIPCC_RTC__
 #endif
 
 #pragma pop_macro("__DEVICE__")

clang/lib/Headers/__clang_hip_runtime_wrapper.h

@@ -18,9 +18,27 @@
 
 #if __HIP__
 
+#if !defined(__HIPCC_RTC__)
 #include <cmath>
 #include <cstdlib>
 #include <stdlib.h>
+#else
+typedef __SIZE_TYPE__ size_t;
+// Define macros which are needed to declare HIP device API's without standard
+// C/C++ headers. This is for readability so that these API's can be written
+// the same way as non-hipRTC use case. These macros need to be popped so that
+// they do not pollute users' name space.
+#pragma push_macro("NULL")
+#pragma push_macro("uint32_t")
+#pragma push_macro("uint64_t")
+#pragma push_macro("CHAR_BIT")
+#pragma push_macro("INT_MAX")
+#define NULL (void *)0
+#define uint32_t __UINT32_TYPE__
+#define uint64_t __UINT64_TYPE__
+#define CHAR_BIT __CHAR_BIT__
+#define INT_MAX __INTMAX_MAX__
+#endif // __HIPCC_RTC__
 
 #define __host__ __attribute__((host))
 #define __device__ __attribute__((device))
@@ -54,6 +72,7 @@ static inline __device__ void *free(void *__ptr) {
 #include <__clang_hip_libdevice_declares.h>
 #include <__clang_hip_math.h>
 
+#if !defined(__HIPCC_RTC__)
 #if !_OPENMP || __HIP_ENABLE_CUDA_WRAPPER_FOR_OPENMP__
 #include <__clang_cuda_math_forward_declares.h>
 #include <__clang_hip_cmath.h>
@@ -62,9 +81,16 @@ static inline __device__ void *free(void *__ptr) {
 #include <algorithm>
 #include <complex>
 #include <new>
+#endif // __HIPCC_RTC__
 #endif // !_OPENMP || __HIP_ENABLE_CUDA_WRAPPER_FOR_OPENMP__
 
 #define __CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__ 1
-
+#if defined(__HIPCC_RTC__)
+#pragma pop_macro("NULL")
+#pragma pop_macro("uint32_t")
+#pragma pop_macro("uint64_t")
+#pragma pop_macro("CHAR_BIT")
+#pragma pop_macro("INT_MAX")
+#endif // __HIPCC_RTC__
 #endif // __HIP__
 #endif // __CLANG_HIP_RUNTIME_WRAPPER_H__

clang/lib/Parse/ParsePragma.cpp

@@ -405,9 +405,11 @@ void Parser::initializePragmaHandlers() {
 
   UnrollHintHandler = std::make_unique<PragmaUnrollHintHandler>("unroll");
   PP.AddPragmaHandler(UnrollHintHandler.get());
+  PP.AddPragmaHandler("GCC", UnrollHintHandler.get());
 
   NoUnrollHintHandler = std::make_unique<PragmaUnrollHintHandler>("nounroll");
   PP.AddPragmaHandler(NoUnrollHintHandler.get());
+  PP.AddPragmaHandler("GCC", NoUnrollHintHandler.get());
 
   UnrollAndJamHintHandler =
       std::make_unique<PragmaUnrollHintHandler>("unroll_and_jam");
@@ -523,9 +525,11 @@ void Parser::resetPragmaHandlers() {
   LoopHintHandler.reset();
 
   PP.RemovePragmaHandler(UnrollHintHandler.get());
+  PP.RemovePragmaHandler("GCC", UnrollHintHandler.get());
   UnrollHintHandler.reset();
 
   PP.RemovePragmaHandler(NoUnrollHintHandler.get());
+  PP.RemovePragmaHandler("GCC", NoUnrollHintHandler.get());
   NoUnrollHintHandler.reset();
 
   PP.RemovePragmaHandler(UnrollAndJamHintHandler.get());

clang/lib/Serialization/ASTReader.cpp

@@ -2760,9 +2760,10 @@ ASTReader::ReadControlBlock(ModuleFile &F,
 
       bool hasErrors = Record[6];
       if (hasErrors && !DisableValidation) {
-        // If requested by the caller, mark modules on error as out-of-date.
-        if (F.Kind == MK_ImplicitModule &&
-            (ClientLoadCapabilities & ARR_TreatModuleWithErrorsAsOutOfDate))
+        // If requested by the caller and the module hasn't already been read
+        // or compiled, mark modules on error as out-of-date.
+        if ((ClientLoadCapabilities & ARR_TreatModuleWithErrorsAsOutOfDate) &&
+            !ModuleMgr.getModuleCache().isPCMFinal(F.FileName))
           return OutOfDate;
 
         if (!AllowASTWithCompilerErrors) {

clang/lib/Tooling/CMakeLists.txt

@@ -24,7 +24,6 @@ string(SUBSTRING ${CMAKE_CURRENT_BINARY_DIR} ${PATH_LIB_END} -1 PATH_TAIL)
 string(CONCAT BINARY_INCLUDE_DIR ${PATH_HEAD} "/include/clang/" ${PATH_TAIL})
 
 if (NOT Python3_EXECUTABLE
-    OR WIN32
     OR APPLE
     OR GENERATOR_IS_MULTI_CONFIG
     OR NOT LLVM_NATIVE_ARCH IN_LIST LLVM_TARGETS_TO_BUILD
@@ -59,6 +58,10 @@ NodeLocationAccessors NodeIntrospection::GetLocations(
     clang::CXXBaseSpecifier const*) {
   return {};
 }
+NodeLocationAccessors NodeIntrospection::GetLocations(
+    clang::TypeLoc const&) {
+  return {};
+}
 NodeLocationAccessors
 NodeIntrospection::GetLocations(clang::DynTypedNode const &) {
   return {};

clang/lib/Tooling/DumpTool/APIData.h

@@ -16,13 +16,11 @@ namespace clang {
 namespace tooling {
 
 struct ClassData {
-
-  bool isEmpty() const {
-    return ASTClassLocations.empty() && ASTClassRanges.empty();
-  }
-
   std::vector<std::string> ASTClassLocations;
   std::vector<std::string> ASTClassRanges;
+  std::vector<std::string> TemplateParms;
+  std::vector<std::string> TypeSourceInfos;
+  std::vector<std::string> TypeLocs;
   // TODO: Extend this with locations available via typelocs etc.
 };