Merge pull request #81 from sx-aurora-dev/feature/merge-upstream-2021…

…0817-2

Feature/merge upstream 20210817 2

clang/include/clang/AST/Type.h

@@ -3456,10 +3456,6 @@ class ConstantMatrixType final : public MatrixType {
 protected:
   friend class ASTContext;
 
-  /// The element type of the matrix.
-  // FIXME: Appears to be unused? There is also MatrixType::ElementType...
-  QualType ElementType;
-
   /// Number of rows and columns.
   unsigned NumRows;
   unsigned NumColumns;

clang/include/clang/Basic/DiagnosticGroups.td

@@ -821,8 +821,10 @@ def ReservedIdentifier : DiagGroup<"reserved-identifier",
 //  under separate flags.
 //
 def UnreachableCodeLoopIncrement : DiagGroup<"unreachable-code-loop-increment">;
+def UnreachableCodeFallthrough : DiagGroup<"unreachable-code-fallthrough">;
 def UnreachableCode : DiagGroup<"unreachable-code",
-                                [UnreachableCodeLoopIncrement]>;
+                                [UnreachableCodeLoopIncrement,
+                                 UnreachableCodeFallthrough]>;
 def UnreachableCodeBreak : DiagGroup<"unreachable-code-break">;
 def UnreachableCodeReturn : DiagGroup<"unreachable-code-return">;
 def UnreachableCodeAggressive : DiagGroup<"unreachable-code-aggressive",

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -682,6 +682,9 @@ def warn_unreachable_return : Warning<
 def warn_unreachable_loop_increment : Warning<
   "loop will run at most once (loop increment never executed)">,
   InGroup<UnreachableCodeLoopIncrement>, DefaultIgnore;
+def warn_unreachable_fallthrough_attr : Warning<
+  "fallthrough annotation in unreachable code">,
+  InGroup<UnreachableCodeFallthrough>, DefaultIgnore;
 def note_unreachable_silence : Note<
   "silence by adding parentheses to mark code as explicitly dead">;
 
@@ -9578,9 +9581,6 @@ def err_fallthrough_attr_outside_switch : Error<
   "fallthrough annotation is outside switch statement">;
 def err_fallthrough_attr_invalid_placement : Error<
   "fallthrough annotation does not directly precede switch label">;
-def warn_fallthrough_attr_unreachable : Warning<
-  "fallthrough annotation in unreachable code">,
-  InGroup<ImplicitFallthrough>, DefaultIgnore;
 
 def warn_unreachable_default : Warning<
   "default label in switch which covers all enumeration values">,

clang/include/clang/Basic/TargetInfo.h

@@ -871,6 +871,11 @@ class TargetInfo : public virtual TransferrableTargetInfo,
   /// across the current set of primary and secondary targets.
   virtual ArrayRef<Builtin::Info> getTargetBuiltins() const = 0;
 
+  /// Returns target-specific min and max values VScale_Range.
+  virtual Optional<std::pair<unsigned, unsigned>>
+  getVScaleRange(const LangOptions &LangOpts) const {
+    return None;
+  }
   /// The __builtin_clz* and __builtin_ctz* built-in
   /// functions are specified to have undefined results for zero inputs, but
   /// on targets that support these operations in a way that provides

clang/lib/AST/ASTDiagnostic.cpp

@@ -1088,6 +1088,9 @@ class TemplateDiff {
             Ty->getAs<TemplateSpecializationType>())
       return TST;
 
+    if (const auto* SubstType = Ty->getAs<SubstTemplateTypeParmType>())
+      Ty = SubstType->getReplacementType();
+
     const RecordType *RT = Ty->getAs<RecordType>();
 
     if (!RT)

clang/lib/Basic/Targets/AArch64.cpp

@@ -424,6 +424,17 @@ ArrayRef<Builtin::Info> AArch64TargetInfo::getTargetBuiltins() const {
                                              Builtin::FirstTSBuiltin);
 }
 
+Optional<std::pair<unsigned, unsigned>>
+AArch64TargetInfo::getVScaleRange(const LangOptions &LangOpts) const {
+  if (LangOpts.ArmSveVectorBits) {
+    unsigned VScale = LangOpts.ArmSveVectorBits / 128;
+    return std::pair<unsigned, unsigned>(VScale, VScale);
+  }
+  if (hasFeature("sve"))
+    return std::pair<unsigned, unsigned>(0, 16);
+  return None;
+}
+
 bool AArch64TargetInfo::hasFeature(StringRef Feature) const {
   return Feature == "aarch64" || Feature == "arm64" || Feature == "arm" ||
          (Feature == "neon" && (FPU & NeonMode)) ||

clang/lib/Basic/Targets/AArch64.h

@@ -96,6 +96,9 @@ class LLVM_LIBRARY_VISIBILITY AArch64TargetInfo : public TargetInfo {
 
   ArrayRef<Builtin::Info> getTargetBuiltins() const override;
 
+  Optional<std::pair<unsigned, unsigned>>
+  getVScaleRange(const LangOptions &LangOpts) const override;
+
   bool hasFeature(StringRef Feature) const override;
   bool handleTargetFeatures(std::vector<std::string> &Features,
                             DiagnosticsEngine &Diags) override;

clang/lib/CodeGen/CodeGenFunction.cpp

@@ -484,11 +484,13 @@ void CodeGenFunction::FinishFunction(SourceLocation EndLoc) {
   //    function.
   CurFn->addFnAttr("min-legal-vector-width", llvm::utostr(LargestVectorWidth));
 
-  // Add vscale attribute if appropriate.
-  if (getLangOpts().ArmSveVectorBits) {
-    unsigned VScale = getLangOpts().ArmSveVectorBits / 128;
-    CurFn->addFnAttr(llvm::Attribute::getWithVScaleRangeArgs(getLLVMContext(),
-                                                             VScale, VScale));
+  // Add vscale_range attribute if appropriate.
+  Optional<std::pair<unsigned, unsigned>> VScaleRange =
+      getContext().getTargetInfo().getVScaleRange(getLangOpts());
+  if (VScaleRange) {
+    CurFn->addFnAttr(llvm::Attribute::getWithVScaleRangeArgs(
+        getLLVMContext(), VScaleRange.getValue().first,
+        VScaleRange.getValue().second));
   }
 
   // If we generated an unreachable return block, delete it now.

clang/lib/Driver/ToolChains/AVR.cpp

@@ -453,11 +453,21 @@ void AVR::Linker::ConstructJob(Compilation &C, const JobAction &JA,
 }
 
 llvm::Optional<std::string> AVRToolChain::findAVRLibcInstallation() const {
+  // Search avr-libc installation according to avr-gcc installation.
+  std::string GCCParent(GCCInstallation.getParentLibPath());
+  std::string Path(GCCParent + "/avr");
+  if (llvm::sys::fs::is_directory(Path))
+    return Path;
+  Path = GCCParent + "/../avr";
+  if (llvm::sys::fs::is_directory(Path))
+    return Path;
+
+  // Search avr-libc installation from possible locations, and return the first
+  // one that exists, if there is no avr-gcc installed.
   for (StringRef PossiblePath : PossibleAVRLibcLocations) {
     std::string Path = getDriver().SysRoot + PossiblePath.str();
-    // Return the first avr-libc installation that exists.
     if (llvm::sys::fs::is_directory(Path))
-      return Optional<std::string>(Path);
+      return Path;
   }
 
   return llvm::None;

clang/lib/Sema/AnalysisBasedWarnings.cpp

@@ -1125,7 +1125,7 @@ namespace {
                 // unreachable in all instantiations of the template.
                 if (!IsTemplateInstantiation)
                   S.Diag(AS->getBeginLoc(),
-                         diag::warn_fallthrough_attr_unreachable);
+                         diag::warn_unreachable_fallthrough_attr);
                 markFallthroughVisited(AS);
                 ++AnnotatedCnt;
                 break;

clang/lib/StaticAnalyzer/Checkers/CheckObjCInstMethSignature.cpp

@@ -6,7 +6,7 @@
 //
 //===----------------------------------------------------------------------===//
 //
-//  This file defines a CheckObjCInstMethSignature, a flow-insenstive check
+//  This file defines a CheckObjCInstMethSignature, a flow-insensitive check
 //  that determines if an Objective-C class interface incorrectly redefines
 //  the method signature in a subclass.
 //

clang/lib/StaticAnalyzer/Checkers/NSErrorChecker.cpp

@@ -6,7 +6,7 @@
 //
 //===----------------------------------------------------------------------===//
 //
-//  This file defines a CheckNSError, a flow-insenstive check
+//  This file defines a CheckNSError, a flow-insensitive check
 //  that determines if an Objective-C class interface correctly returns
 //  a non-void return type.
 //

clang/lib/StaticAnalyzer/Core/HTMLDiagnostics.cpp

@@ -1289,15 +1289,32 @@ var findNum = function() {
     return out;
 };
 
+var classListAdd = function(el, theClass) {
+  if(!el.className.baseVal)
+    el.className += " " + theClass;
+  else
+    el.className.baseVal += " " + theClass;
+};
+
+var classListRemove = function(el, theClass) {
+  var className = (!el.className.baseVal) ?
+      el.className : el.className.baseVal;
+    className = className.replace(" " + theClass, "");
+  if(!el.className.baseVal)
+    el.className = className;
+  else
+    el.className.baseVal = className;
+};
+
 var scrollTo = function(el) {
     querySelectorAllArray(".selected").forEach(function(s) {
-        s.classList.remove("selected");
+      classListRemove(s, "selected");
     });
-    el.classList.add("selected");
+    classListAdd(el, "selected");
     window.scrollBy(0, el.getBoundingClientRect().top -
         (window.innerHeight / 2));
     highlightArrowsForSelectedEvent();
-}
+};
 
 var move = function(num, up, numItems) {
   if (num == 1 && up || num == numItems - 1 && !up) {
@@ -1332,9 +1349,11 @@ window.addEventListener("keydown", function (event) {
   if (event.defaultPrevented) {
     return;
   }
-  if (event.key == "j") {
+  // key 'j'
+  if (event.keyCode == 74) {
     navigateTo(/*up=*/false);
-  } else if (event.key == "k") {
+  // key 'k'
+  } else if (event.keyCode == 75) {
     navigateTo(/*up=*/true);
   } else {
     return;
@@ -1350,8 +1369,11 @@ StringRef HTMLDiagnostics::generateArrowDrawingJavascript() {
 <script type='text/javascript'>
 // Return range of numbers from a range [lower, upper).
 function range(lower, upper) {
-  const size = upper - lower;
-  return Array.from(new Array(size), (x, i) => i + lower);
+  var array = [];
+  for (var i = lower; i <= upper; ++i) {
+      array.push(i);
+  }
+  return array;
 }
 
 var getRelatedArrowIndices = function(pathId) {
@@ -1371,7 +1393,9 @@ var highlightArrowsForSelectedEvent = function() {
   const arrowIndicesToHighlight = getRelatedArrowIndices(selectedNum);
   arrowIndicesToHighlight.forEach((index) => {
     var arrow = document.querySelector("#arrow" + index);
-    arrow.classList.add("selected");
+    if(arrow) {
+      classListAdd(arrow, "selected")
+    }
   });
 }
 

clang/test/CodeGen/arm-sve-vector-bits-vscale-range.c

@@ -3,10 +3,13 @@
 // RUN: %clang_cc1 -triple aarch64-none-linux-gnu -target-feature +sve -msve-vector-bits=512 -S -emit-llvm -o - %s | FileCheck %s -D#VBITS=512
 // RUN: %clang_cc1 -triple aarch64-none-linux-gnu -target-feature +sve -msve-vector-bits=1024 -S -emit-llvm -o - %s | FileCheck %s -D#VBITS=1024
 // RUN: %clang_cc1 -triple aarch64-none-linux-gnu -target-feature +sve -msve-vector-bits=2048 -S -emit-llvm -o - %s | FileCheck %s -D#VBITS=2048
+// RUN: %clang_cc1 -triple aarch64-none-linux-gnu -target-feature +sve2 -msve-vector-bits=128 -S -emit-llvm -o - %s | FileCheck %s -D#VBITS=128
+// RUN: %clang_cc1 -triple aarch64-none-linux-gnu -target-feature +sve2 -msve-vector-bits=256 -S -emit-llvm -o - %s | FileCheck %s -D#VBITS=256
+// RUN: %clang_cc1 -triple aarch64-none-linux-gnu -target-feature +sve2 -msve-vector-bits=scalable -S -emit-llvm -o - %s | FileCheck %s --check-prefix=CHECK-NONE
 // RUN: %clang_cc1 -triple aarch64-none-linux-gnu -target-feature +sve -msve-vector-bits=scalable -S -emit-llvm -o - %s | FileCheck %s --check-prefix=CHECK-NONE
 // RUN: %clang_cc1 -triple aarch64-none-linux-gnu -target-feature +sve -S -emit-llvm -o - %s | FileCheck %s --check-prefix=CHECK-NONE
 
 // CHECK-LABEL: @func() #0
 // CHECK: attributes #0 = { {{.*}} vscale_range([[#div(VBITS,128)]],[[#div(VBITS,128)]]) {{.*}} }
-// CHECK-NONE-NOT: vscale_range
+// CHECK-NONE: attributes #0 = { {{.*}} vscale_range(0,16) {{.*}} }
 void func() {}

clang/test/CodeGen/thinlto-distributed-newpm.ll

@@ -134,7 +134,6 @@
 ; CHECK-O: Running pass: LoopSimplifyPass on main
 ; CHECK-O: Running analysis: LoopAnalysis on main
 ; CHECK-O: Running pass: LCSSAPass on main
-; CHECK-O: Running analysis: MemorySSAAnalysis on main
 ; CHECK-O: Running analysis: AAManager on main
 ; CHECK-O: Running analysis: BasicAA on main
 ; CHECK-O: Running analysis: ScalarEvolutionAnalysis on main

clang/test/CodeGenCUDA/atomics-remarks-gfx90a.cu

@@ -0,0 +1,16 @@
+// RUN: %clang_cc1 %s -triple=amdgcn-amd-amdhsa -fcuda-is-device \
+// RUN:   -target-cpu gfx90a -Rpass=atomic-expand -S -o - 2>&1 | \
+// RUN:   FileCheck %s --check-prefix=GFX90A-CAS
+
+// REQUIRES: amdgpu-registered-target
+
+#include "Inputs/cuda.h"
+#include <stdatomic.h>
+
+// GFX90A-CAS: A compare and swap loop was generated for an atomic fadd operation at system memory scope
+// GFX90A-CAS-LABEL: _Z14atomic_add_casPf
+// GFX90A-CAS:  flat_atomic_cmpswap v0, v[2:3], v[4:5] glc
+// GFX90A-CAS:  s_cbranch_execnz
+__device__ float atomic_add_cas(float *p) {
+  return __atomic_fetch_add(p, 1.0f, memory_order_relaxed);
+}

clang/test/CodeGenOpenCL/atomics-remarks-gfx90a.cl

@@ -0,0 +1,43 @@
+// RUN: %clang_cc1 %s -cl-std=CL2.0 -O0 -triple=amdgcn-amd-amdhsa -target-cpu gfx90a \
+// RUN:     -Rpass=atomic-expand -S -o - 2>&1 | \
+// RUN:     FileCheck %s --check-prefix=REMARK
+
+// RUN: %clang_cc1 %s -cl-std=CL2.0 -O0 -triple=amdgcn-amd-amdhsa -target-cpu gfx90a \
+// RUN:     -Rpass=atomic-expand -S -emit-llvm -o - 2>&1 | \
+// RUN:     FileCheck %s --check-prefix=GFX90A-CAS
+
+// REQUIRES: amdgpu-registered-target
+
+typedef enum memory_order {
+  memory_order_relaxed = __ATOMIC_RELAXED,
+  memory_order_acquire = __ATOMIC_ACQUIRE,
+  memory_order_release = __ATOMIC_RELEASE,
+  memory_order_acq_rel = __ATOMIC_ACQ_REL,
+  memory_order_seq_cst = __ATOMIC_SEQ_CST
+} memory_order;
+
+typedef enum memory_scope {
+  memory_scope_work_item = __OPENCL_MEMORY_SCOPE_WORK_ITEM,
+  memory_scope_work_group = __OPENCL_MEMORY_SCOPE_WORK_GROUP,
+  memory_scope_device = __OPENCL_MEMORY_SCOPE_DEVICE,
+  memory_scope_all_svm_devices = __OPENCL_MEMORY_SCOPE_ALL_SVM_DEVICES,
+#if defined(cl_intel_subgroups) || defined(cl_khr_subgroups)
+  memory_scope_sub_group = __OPENCL_MEMORY_SCOPE_SUB_GROUP
+#endif
+} memory_scope;
+
+// REMARK: remark: A compare and swap loop was generated for an atomic fadd operation at workgroup-one-as memory scope [-Rpass=atomic-expand]
+// REMARK: remark: A compare and swap loop was generated for an atomic fadd operation at agent-one-as memory scope [-Rpass=atomic-expand]
+// REMARK: remark: A compare and swap loop was generated for an atomic fadd operation at one-as memory scope [-Rpass=atomic-expand]
+// REMARK: remark: A compare and swap loop was generated for an atomic fadd operation at wavefront-one-as memory scope [-Rpass=atomic-expand]
+// GFX90A-CAS-LABEL: @atomic_cas
+// GFX90A-CAS: atomicrmw fadd float addrspace(1)* {{.*}} syncscope("workgroup-one-as") monotonic
+// GFX90A-CAS: atomicrmw fadd float addrspace(1)* {{.*}} syncscope("agent-one-as") monotonic
+// GFX90A-CAS: atomicrmw fadd float addrspace(1)* {{.*}} syncscope("one-as") monotonic
+// GFX90A-CAS: atomicrmw fadd float addrspace(1)* {{.*}} syncscope("wavefront-one-as") monotonic
+float atomic_cas(__global atomic_float *d, float a) {
+  float ret1 = __opencl_atomic_fetch_add(d, a, memory_order_relaxed, memory_scope_work_group);
+  float ret2 = __opencl_atomic_fetch_add(d, a, memory_order_relaxed, memory_scope_device);
+  float ret3 = __opencl_atomic_fetch_add(d, a, memory_order_relaxed, memory_scope_all_svm_devices);
+  float ret4 = __opencl_atomic_fetch_add(d, a, memory_order_relaxed, memory_scope_sub_group);
+}

clang/test/Driver/Inputs/empty-elf-template.yaml

@@ -0,0 +1,5 @@
+--- !ELF
+FileHeader:
+  Class: ELFCLASS[[BITS]]
+  Data:  ELFDATA2[[ENCODING]]
+  Type:  ET_REL

clang/test/Driver/avr-toolchain.c

@@ -6,12 +6,24 @@
 // CHECK1-SAME: "-resource-dir" "[[RESOURCE:[^"]+]]"
 // CHECK1-SAME: "-isysroot" "[[SYSROOT:[^"]+/basic_avr_tree]]"
 // CHECK1-SAME: "-internal-isystem"
-// CHECK1-SAME: {{^}} "[[SYSROOT]]/usr/lib/avr/include"
+// CHECK1-SAME: {{^}} "[[SYSROOT]]/usr/lib/gcc/avr/5.4.0/../../../avr/include"
 // CHECK1-NOT:  "-L
 // CHECK1:      avr-ld"
 // CHECK1-SAME: "-o" "a.out"
 // CHECK1-SAME: {{^}} "--gc-sections"
 
+// RUN: %clang %s -### -target avr --sysroot %S/Inputs/basic_avr_tree_2/opt/local -S 2>&1 | FileCheck --check-prefix=CHECK2 %s
+// CHECK2: clang{{.*}} "-cc1" "-triple" "avr"
+// CHECK2-SAME: "-isysroot" "[[SYSROOT:[^"]+/basic_avr_tree_2/opt/local]]"
+// CHECK2-SAME: "-internal-isystem"
+// CHECK2-SAME: {{^}} "[[SYSROOT]]/lib/gcc/avr/10.3.0/../../../../avr/include"
+
+// RUN: %clang %s -### -target avr --sysroot %S/Inputs/basic_avr_tree_2 -S 2>&1 | FileCheck --check-prefix=CHECK3 %s
+// CHECK3: clang{{.*}} "-cc1" "-triple" "avr"
+// CHECK3-SAME: "-isysroot" "[[SYSROOT:[^"]+/basic_avr_tree_2]]"
+// CHECK3-SAME: "-internal-isystem"
+// CHECK3-SAME: {{^}} "[[SYSROOT]]/usr/avr/include"
+
 // RUN: %clang %s -### -target avr 2>&1 | FileCheck -check-prefix=CC1 %s
 // CC1: clang{{.*}} "-cc1" "-triple" "avr" {{.*}} "-fno-use-init-array"
 

clang/test/Driver/clang-offload-wrapper.c

@@ -19,9 +19,10 @@
 //
 // Check bitcode produced by the wrapper tool.
 //
-// RUN: clang-offload-wrapper -target=x86_64-pc-linux-gnu -o %t.wrapper.bc %t.tgt
+// RUN: clang-offload-wrapper -target=x86_64-pc-linux-gnu -o %t.wrapper.bc %t.tgt 2>&1 | FileCheck %s --check-prefix ELF-WARNING
 // RUN: llvm-dis %t.wrapper.bc -o - | FileCheck %s --check-prefix CHECK-IR
 
+// ELF-WARNING: is not an ELF image, so notes cannot be added to it.
 // CHECK-IR: target triple = "x86_64-pc-linux-gnu"
 
 // CHECK-IR-DAG: [[ENTTY:%.+]] = type { i8*, i8*, i{{32|64}}, i32, i32 }
@@ -53,3 +54,24 @@
 // CHECK-IR:   ret void
 
 // CHECK-IR: declare void @__tgt_unregister_lib([[DESCTY]]*)
+
+// Check that clang-offload-wrapper adds LLVMOMPOFFLOAD notes
+// into the ELF offload images:
+// RUN: yaml2obj %S/Inputs/empty-elf-template.yaml -o %t.64le -DBITS=64 -DENCODING=LSB
+// RUN: clang-offload-wrapper -target=x86_64-pc-linux-gnu -o %t.wrapper.elf64le.bc %t.64le
+// RUN: llvm-dis %t.wrapper.elf64le.bc -o - | FileCheck %s --check-prefix OMPNOTES
+// RUN: yaml2obj %S/Inputs/empty-elf-template.yaml -o %t.64be -DBITS=64 -DENCODING=MSB
+// RUN: clang-offload-wrapper -target=x86_64-pc-linux-gnu -o %t.wrapper.elf64be.bc %t.64be
+// RUN: llvm-dis %t.wrapper.elf64be.bc -o - | FileCheck %s --check-prefix OMPNOTES
+// RUN: yaml2obj %S/Inputs/empty-elf-template.yaml -o %t.32le -DBITS=32 -DENCODING=LSB
+// RUN: clang-offload-wrapper -target=x86_64-pc-linux-gnu -o %t.wrapper.elf32le.bc %t.32le
+// RUN: llvm-dis %t.wrapper.elf32le.bc -o - | FileCheck %s --check-prefix OMPNOTES
+// RUN: yaml2obj %S/Inputs/empty-elf-template.yaml -o %t.32be -DBITS=32 -DENCODING=MSB
+// RUN: clang-offload-wrapper -target=x86_64-pc-linux-gnu -o %t.wrapper.elf32be.bc %t.32be
+// RUN: llvm-dis %t.wrapper.elf32be.bc -o - | FileCheck %s --check-prefix OMPNOTES
+
+// There is no clean way for extracting the offload image
+// from the object file currently, so try to find
+// the inserted ELF notes in the device image variable's
+// initializer:
+// OMPNOTES: @{{.+}} = internal unnamed_addr constant [{{[0-9]+}} x i8] c"{{.*}}LLVMOMPOFFLOAD{{.*}}LLVMOMPOFFLOAD{{.*}}LLVMOMPOFFLOAD{{.*}}"