Merged master:80cd6b6e043 into amd-gfx:28ca29a47cf

Local branch amd-gfx 28ca29a Merged master:6b2431e0c2a into amd-gfx:fb843514f54 Remote branch master 80cd6b6 [SLP] Enhance SLPVectorizer to vectorize vector aggregate
jaebaek · Nov 22, 2019 · 148e72f · 148e72f
2 parents 28ca29a + 80cd6b6
commit 148e72f
Show file tree

Hide file tree

Showing 41 changed files with 823 additions and 108 deletions.
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -11138,33 +11138,22 @@ bool checkContext<OMP_CTX_SET_device, OMP_CTX_kind, CodeGenModule &>(
 bool matchesContext(CodeGenModule &CGM,
                     const CompleteOMPContextSelectorData &ContextData) {
   for (const OMPContextSelectorData &Data : ContextData) {
-    switch (Data.CtxSet) {
-    case OMP_CTX_SET_implementation:
-      switch (Data.Ctx) {
-      case OMP_CTX_vendor:
-        if (!checkContext<OMP_CTX_SET_implementation, OMP_CTX_vendor>(Data))
-          return false;
-        break;
-      case OMP_CTX_kind:
-      case OMP_CTX_unknown:
-        llvm_unreachable(
-            "Unexpected context selector kind in implementation set.");
-      }
+    switch (Data.Ctx) {
+    case OMP_CTX_vendor:
+      assert(Data.CtxSet == OMP_CTX_SET_implementation &&
+             "Expected implementation context selector set.");
+      if (!checkContext<OMP_CTX_SET_implementation, OMP_CTX_vendor>(Data))
+        return false;
       break;
-    case OMP_CTX_SET_device:
-      switch (Data.Ctx) {
-      case OMP_CTX_kind:
-        if (!checkContext<OMP_CTX_SET_device, OMP_CTX_kind, CodeGenModule &>(
-                Data, CGM))
-          return false;
-        break;
-      case OMP_CTX_vendor:
-      case OMP_CTX_unknown:
-        llvm_unreachable("Unexpected context selector kind in device set.");
-      }
+    case OMP_CTX_kind:
+      assert(Data.CtxSet == OMP_CTX_SET_device &&
+             "Expected device context selector set.");
+      if (!checkContext<OMP_CTX_SET_device, OMP_CTX_kind, CodeGenModule &>(Data,
+                                                                           CGM))
+        return false;
       break;
-    case OMP_CTX_SET_unknown:
-      llvm_unreachable("Unexpected context selector set kind.");
+    case OMP_CTX_unknown:
+      llvm_unreachable("Unknown context selector kind.");
     }
   }
   return true;
@@ -11184,32 +11173,21 @@ translateAttrToContextSelectorData(ASTContext &C,
     Data.back().Ctx = Ctx;
     const Expr *Score = *std::next(A->scores_begin(), I);
     Data.back().Score = Score->EvaluateKnownConstInt(C);
-    switch (CtxSet) {
-    case OMP_CTX_SET_implementation:
-      switch (Ctx) {
-      case OMP_CTX_vendor:
-        Data.back().Names =
-            llvm::makeArrayRef(A->implVendors_begin(), A->implVendors_end());
-        break;
-      case OMP_CTX_kind:
-      case OMP_CTX_unknown:
-        llvm_unreachable(
-            "Unexpected context selector kind in implementation set.");
-      }
+    switch (Ctx) {
+    case OMP_CTX_vendor:
+      assert(CtxSet == OMP_CTX_SET_implementation &&
+             "Expected implementation context selector set.");
+      Data.back().Names =
+          llvm::makeArrayRef(A->implVendors_begin(), A->implVendors_end());
       break;
-    case OMP_CTX_SET_device:
-      switch (Ctx) {
-      case OMP_CTX_kind:
-        Data.back().Names =
-            llvm::makeArrayRef(A->deviceKinds_begin(), A->deviceKinds_end());
-        break;
-      case OMP_CTX_vendor:
-      case OMP_CTX_unknown:
-        llvm_unreachable("Unexpected context selector kind in device set.");
-      }
+    case OMP_CTX_kind:
+      assert(CtxSet == OMP_CTX_SET_device &&
+             "Expected device context selector set.");
+      Data.back().Names =
+          llvm::makeArrayRef(A->deviceKinds_begin(), A->deviceKinds_end());
       break;
-    case OMP_CTX_SET_unknown:
-      llvm_unreachable("Unexpected context selector set kind.");
+    case OMP_CTX_unknown:
+      llvm_unreachable("Unknown context selector kind.");
     }
   }
   return Data;

diff --git a/clang/lib/Sema/SemaCoroutine.cpp b/clang/lib/Sema/SemaCoroutine.cpp
@@ -1527,8 +1527,8 @@ bool Sema::buildCoroutineParameterMoves(SourceLocation Loc) {
   auto *FD = cast<FunctionDecl>(CurContext);
 
   auto *ScopeInfo = getCurFunction();
-  assert(ScopeInfo->CoroutineParameterMoves.empty() &&
-         "Should not build parameter moves twice");
+  if (!ScopeInfo->CoroutineParameterMoves.empty())
+    return false;
 
   for (auto *PD : FD->parameters()) {
     if (PD->getType()->isDependentType())

diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp
@@ -1763,7 +1763,12 @@ void Sema::checkOpenMPDeviceExpr(const Expr *E) {
 }
 
 static OpenMPDefaultmapClauseKind
-getVariableCategoryFromDecl(const ValueDecl *VD) {
+getVariableCategoryFromDecl(const LangOptions &LO, const ValueDecl *VD) {
+  if (LO.OpenMP <= 45) {
+    if (VD->getType().getNonReferenceType()->isScalarType())
+      return OMPC_DEFAULTMAP_scalar;
+    return OMPC_DEFAULTMAP_aggregate;
+  }
   if (VD->getType().getNonReferenceType()->isAnyPointerType())
     return OMPC_DEFAULTMAP_pointer;
   if (VD->getType().getNonReferenceType()->isScalarType())
@@ -1894,8 +1899,8 @@ bool Sema::isOpenMPCapturedByRef(const ValueDecl *D, unsigned Level,
           (DSAStack->isForceCaptureByReferenceInTargetExecutable() &&
            !Ty->isAnyPointerType()) ||
           !Ty->isScalarType() ||
-          DSAStack->isDefaultmapCapturedByRef(Level,
-                                              getVariableCategoryFromDecl(D)) ||
+          DSAStack->isDefaultmapCapturedByRef(
+              Level, getVariableCategoryFromDecl(LangOpts, D)) ||
           DSAStack->hasExplicitDSA(
               D, [](OpenMPClauseKind K) { return K == OMPC_reduction; }, Level);
     }
@@ -2112,8 +2117,8 @@ void Sema::setOpenMPCaptureKind(FieldDecl *FD, const ValueDecl *D,
     if (DSAStack->hasExplicitDirective(isOpenMPTargetExecutionDirective,
                                        NewLevel)) {
       OMPC = OMPC_map;
-      if (DSAStack->mustBeFirstprivateAtLevel(NewLevel,
-                                              getVariableCategoryFromDecl(D)))
+      if (DSAStack->mustBeFirstprivateAtLevel(
+              NewLevel, getVariableCategoryFromDecl(LangOpts, D)))
         OMPC = OMPC_firstprivate;
       break;
     }
@@ -2944,7 +2949,8 @@ class DSAAttrChecker final : public StmtVisitor<DSAAttrChecker, void> {
       // data-haring attribute clause (including a data-sharing attribute
       // clause on a combined construct where target. is one of the
       // constituent constructs), or an is_device_ptr clause.
-      OpenMPDefaultmapClauseKind ClauseKind = getVariableCategoryFromDecl(VD);
+      OpenMPDefaultmapClauseKind ClauseKind =
+          getVariableCategoryFromDecl(SemaRef.getLangOpts(), VD);
       if (SemaRef.getLangOpts().OpenMP >= 50) {
         bool IsModifierNone = Stack->getDefaultmapModifier(ClauseKind) ==
                               OMPC_DEFAULTMAP_MODIFIER_none;
@@ -5398,29 +5404,19 @@ void Sema::ActOnOpenMPDeclareVariantDirective(
       else
         Score = ActOnIntegerConstant(SourceLocation(), 0).get();
     }
-    switch (CtxSet) {
-    case OMP_CTX_SET_implementation:
-      switch (Ctx) {
-      case OMP_CTX_vendor:
-        ImplVendors.append(D.Names.begin(), D.Names.end());
-        break;
-      case OMP_CTX_kind:
-      case OMP_CTX_unknown:
-        llvm_unreachable("Unexpected context selector kind.");
-      }
+    switch (Ctx) {
+    case OMP_CTX_vendor:
+      assert(CtxSet == OMP_CTX_SET_implementation &&
+             "Expected implementation context selector set.");
+      ImplVendors.append(D.Names.begin(), D.Names.end());
       break;
-    case OMP_CTX_SET_device:
-      switch (Ctx) {
-      case OMP_CTX_kind:
-        DeviceKinds.append(D.Names.begin(), D.Names.end());
-        break;
-      case OMP_CTX_vendor:
-      case OMP_CTX_unknown:
-        llvm_unreachable("Unexpected context selector kind.");
-      }
+    case OMP_CTX_kind:
+      assert(CtxSet == OMP_CTX_SET_device &&
+             "Expected device context selector set.");
+      DeviceKinds.append(D.Names.begin(), D.Names.end());
       break;
-    case OMP_CTX_SET_unknown:
-      llvm_unreachable("Unexpected context selector set kind.");
+    case OMP_CTX_unknown:
+      llvm_unreachable("Unknown context selector kind.");
     }
     IsError = IsError || !Score;
     CtxSets.push_back(CtxSet);

diff --git a/clang/test/OpenMP/target_map_codegen.cpp b/clang/test/OpenMP/target_map_codegen.cpp
@@ -656,15 +656,15 @@ void implicit_maps_pointer (){
 
 // CK11-LABEL: @.__omp_offloading_{{.*}}implicit_maps_double_complex{{.*}}_l678.region_id = weak constant i8 0
 
-// CK11-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]
+// CK11-DAG: [[SIZES:@.+]] = {{.+}}constant [2 x i64] [i64 16, i64 {{8|4}}]
 // Map types: OMP_MAP_TO  | OMP_MAP_FROM | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 547
-// CK11-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 547]
+// CK11-DAG: [[TYPES:@.+]] = {{.+}}constant [2 x i64] [i64 547, i64 547]
 
 // CK11-LABEL: implicit_maps_double_complex{{.*}}(
-void implicit_maps_double_complex (int a){
+void implicit_maps_double_complex (int a, int *b){
   double _Complex dc = (double)a;
 
-  // CK11-DAG: call i32 @__tgt_target(i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}})
+  // CK11-DAG: call i32 @__tgt_target(i64 {{.+}}, i8* {{.+}}, i32 2, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}})
   // CK11-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
   // CK11-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
   // CK11-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
@@ -674,14 +674,14 @@ void implicit_maps_double_complex (int a){
   // CK11-DAG: store { double, double }* [[PTR:%[^,]+]], { double, double }** [[CBP1]]
   // CK11-DAG: store { double, double }* [[PTR]], { double, double }** [[CP1]]
 
-  // CK11: call void [[KERNEL:@.+]]({ double, double }* [[PTR]])
+  // CK11: call void [[KERNEL:@.+]]({ double, double }* [[PTR]], i32** %{{.+}})
   #pragma omp target defaultmap(tofrom:scalar)
   {
-   dc *= dc;
+   dc *= dc; *b = 1;
   }
 }
 
-// CK11: define internal void [[KERNEL]]({ double, double }* {{.*}}[[ARG:%.+]])
+// CK11: define internal void [[KERNEL]]({ double, double }* {{.*}}[[ARG:%.+]], i32** {{.*}})
 // CK11: [[ADDR:%.+]] = alloca { double, double }*,
 // CK11: store { double, double }* [[ARG]], { double, double }** [[ADDR]],
 // CK11: [[REF:%.+]] = load { double, double }*, { double, double }** [[ADDR]],

diff --git a/clang/test/SemaCXX/coroutines.cpp b/clang/test/SemaCXX/coroutines.cpp
@@ -87,6 +87,11 @@ int no_promise_type() { // expected-error {{this function cannot be a coroutine:
   co_await a;
 }
 
+int no_promise_type_multiple_awaits(int) { // expected-error {{this function cannot be a coroutine: 'std::experimental::coroutine_traits<int, int>' has no member named 'promise_type'}}
+  co_await a;
+  co_await a;
+}
+
 template <>
 struct std::experimental::coroutine_traits<double, double> { typedef int promise_type; };
 double bad_promise_type(double) { // expected-error {{this function cannot be a coroutine: 'experimental::coroutine_traits<double, double>::promise_type' (aka 'int') is not a class}}

diff --git a/clang/tools/clang-scan-deps/ClangScanDeps.cpp b/clang/tools/clang-scan-deps/ClangScanDeps.cpp
@@ -13,6 +13,7 @@
 #include "clang/Tooling/DependencyScanning/DependencyScanningWorker.h"
 #include "clang/Tooling/JSONCompilationDatabase.h"
 #include "llvm/Support/CommandLine.h"
+#include "llvm/Support/FileUtilities.h"
 #include "llvm/Support/InitLLVM.h"
 #include "llvm/Support/Program.h"
 #include "llvm/Support/Signals.h"
@@ -39,6 +40,64 @@ class SharedStream {
   raw_ostream &OS;
 };
 
+class ResourceDirectoryCache {
+public:
+  /// findResourceDir finds the resource directory relative to the clang
+  /// compiler being used in Args, by running it with "-print-resource-dir"
+  /// option and cache the results for reuse. \returns resource directory path
+  /// associated with the given invocation command or empty string if the
+  /// compiler path is NOT an absolute path.
+  StringRef findResourceDir(const tooling::CommandLineArguments &Args) {
+    if (Args.size() < 1)
+      return "";
+
+    const std::string &ClangBinaryPath = Args[0];
+    if (!llvm::sys::path::is_absolute(ClangBinaryPath))
+      return "";
+
+    const std::string &ClangBinaryName =
+        llvm::sys::path::filename(ClangBinaryPath);
+
+    std::unique_lock<std::mutex> LockGuard(CacheLock);
+    const auto &CachedResourceDir = Cache.find(ClangBinaryPath);
+    if (CachedResourceDir != Cache.end())
+      return CachedResourceDir->second;
+
+    std::vector<StringRef> PrintResourceDirArgs{ClangBinaryName,
+                                                "-print-resource-dir"};
+    llvm::SmallString<64> OutputFile, ErrorFile;
+    llvm::sys::fs::createTemporaryFile("print-resource-dir-output",
+                                       "" /*no-suffix*/, OutputFile);
+    llvm::sys::fs::createTemporaryFile("print-resource-dir-error",
+                                       "" /*no-suffix*/, ErrorFile);
+    llvm::FileRemover OutputRemover(OutputFile.c_str());
+    llvm::FileRemover ErrorRemover(ErrorFile.c_str());
+    llvm::Optional<StringRef> Redirects[] = {
+        {""}, // Stdin
+        StringRef(OutputFile),
+        StringRef(ErrorFile),
+    };
+    if (const int RC = llvm::sys::ExecuteAndWait(
+            ClangBinaryPath, PrintResourceDirArgs, {}, Redirects)) {
+      auto ErrorBuf = llvm::MemoryBuffer::getFile(ErrorFile.c_str());
+      llvm::errs() << ErrorBuf.get()->getBuffer();
+      return "";
+    }
+
+    auto OutputBuf = llvm::MemoryBuffer::getFile(OutputFile.c_str());
+    if (!OutputBuf)
+      return "";
+    StringRef Output = OutputBuf.get()->getBuffer().rtrim('\n');
+
+    Cache[ClangBinaryPath] = Output.str();
+    return Cache[ClangBinaryPath];
+  }
+
+private:
+  std::map<std::string, std::string> Cache;
+  std::mutex CacheLock;
+};
+
 llvm::cl::opt<bool> Help("h", llvm::cl::desc("Alias for -help"),
                          llvm::cl::Hidden);
 
@@ -169,12 +228,15 @@ int main(int argc, const char **argv) {
   auto AdjustingCompilations =
       std::make_unique<tooling::ArgumentsAdjustingCompilations>(
           std::move(Compilations));
+  ResourceDirectoryCache ResourceDirCache;
   AdjustingCompilations->appendArgumentsAdjuster(
-      [](const tooling::CommandLineArguments &Args, StringRef FileName) {
+      [&ResourceDirCache](const tooling::CommandLineArguments &Args,
+                          StringRef FileName) {
         std::string LastO = "";
         bool HasMT = false;
         bool HasMQ = false;
         bool HasMD = false;
+        bool HasResourceDir = false;
         // We need to find the last -o value.
         if (!Args.empty()) {
           std::size_t Idx = Args.size() - 1;
@@ -188,6 +250,8 @@ int main(int argc, const char **argv) {
                 HasMQ = true;
               if (Args[Idx] == "-MD")
                 HasMD = true;
+              if (Args[Idx] == "-resource-dir")
+                HasResourceDir = true;
             }
             --Idx;
           }
@@ -215,6 +279,15 @@ int main(int argc, const char **argv) {
         AdjustedArgs.push_back("-Xclang");
         AdjustedArgs.push_back("-sys-header-deps");
         AdjustedArgs.push_back("-Wno-error");
+
+        if (!HasResourceDir) {
+          StringRef ResourceDir =
+              ResourceDirCache.findResourceDir(Args);
+          if (!ResourceDir.empty()) {
+            AdjustedArgs.push_back("-resource-dir");
+            AdjustedArgs.push_back(ResourceDir);
+          }
+        }
         return AdjustedArgs;
       });
   AdjustingCompilations->appendArgumentsAdjuster(

diff --git a/llvm/lib/MC/XCOFFObjectWriter.cpp b/llvm/lib/MC/XCOFFObjectWriter.cpp
@@ -150,6 +150,7 @@ class XCOFFObjectWriter : public MCObjectWriter {
   // the sections. Should have one for each set of csects that get mapped into
   // the same section and get handled in a 'similar' way.
   CsectGroup ProgramCodeCsects;
+  CsectGroup ReadOnlyCsects;
   CsectGroup DataCsects;
   CsectGroup FuncDSCsects;
   CsectGroup TOCCsects;
@@ -219,7 +220,7 @@ XCOFFObjectWriter::XCOFFObjectWriter(
     : W(OS, support::big), TargetObjectWriter(std::move(MOTW)),
       Strings(StringTableBuilder::XCOFF),
       Text(".text", XCOFF::STYP_TEXT, /* IsVirtual */ false,
-           CsectGroups{&ProgramCodeCsects}),
+           CsectGroups{&ProgramCodeCsects, &ReadOnlyCsects}),
       Data(".data", XCOFF::STYP_DATA, /* IsVirtual */ false,
            CsectGroups{&DataCsects, &FuncDSCsects, &TOCCsects}),
       BSS(".bss", XCOFF::STYP_BSS, /* IsVirtual */ true,
@@ -245,6 +246,10 @@ CsectGroup &XCOFFObjectWriter::getCsectGroup(const MCSectionXCOFF *MCSec) {
     assert(XCOFF::XTY_SD == MCSec->getCSectType() &&
            "Only an initialized csect can contain program code.");
     return ProgramCodeCsects;
+  case XCOFF::XMC_RO:
+    assert(XCOFF::XTY_SD == MCSec->getCSectType() &&
+           "Only an initialized csect can contain read only data.");
+    return ReadOnlyCsects;
   case XCOFF::XMC_RW:
     if (XCOFF::XTY_CM == MCSec->getCSectType())
       return BSSCsects;