Merge branch 'main' into fix-ProTypeUnionAccessCheck

Author: nicovank

clang/include/clang/Basic/AttrDocs.td

@@ -2669,7 +2669,7 @@ sign. For example:
 
   .. code-block:: c++
 
-    __attribute__((target_clones("sha2+memtag2", "fcma+sve2-pmull128")))
+    __attribute__((target_clones("sha2+memtag", "fcma+sve2-pmull128")))
     void foo() {}
 
 For every multiversioned function a ``default`` (fallback) implementation

clang/include/clang/Basic/BuiltinsAMDGPU.def

@@ -224,8 +224,8 @@ TARGET_BUILTIN(__builtin_amdgcn_frexp_exph, "sh", "nc", "16-bit-insts")
 TARGET_BUILTIN(__builtin_amdgcn_fracth, "hh", "nc", "16-bit-insts")
 TARGET_BUILTIN(__builtin_amdgcn_classh, "bhi", "nc", "16-bit-insts")
 TARGET_BUILTIN(__builtin_amdgcn_s_memrealtime, "WUi", "n", "s-memrealtime")
-TARGET_BUILTIN(__builtin_amdgcn_mov_dpp, "iiIiIiIiIb", "nc", "dpp")
-TARGET_BUILTIN(__builtin_amdgcn_update_dpp, "iiiIiIiIiIb", "nc", "dpp")
+TARGET_BUILTIN(__builtin_amdgcn_mov_dpp, "iiIiIiIiIb", "nct", "dpp")
+TARGET_BUILTIN(__builtin_amdgcn_update_dpp, "iiiIiIiIiIb", "nct", "dpp")
 TARGET_BUILTIN(__builtin_amdgcn_s_dcache_wb, "v", "n", "gfx8-insts")
 TARGET_BUILTIN(__builtin_amdgcn_perm, "UiUiUiUi", "nc", "gfx8-insts")
 

clang/lib/Basic/Targets/AArch64.cpp

@@ -784,7 +784,7 @@ bool AArch64TargetInfo::hasFeature(StringRef Feature) const {
       .Case("sme-fa64", HasSMEFA64)
       .Case("sme-f16f16", HasSMEF16F16)
       .Case("sme-b16b16", HasSMEB16B16)
-      .Cases("memtag", "memtag2", HasMTE)
+      .Case("memtag", HasMTE)
       .Case("sb", HasSB)
       .Case("predres", HasPredRes)
       .Cases("ssbs", "ssbs2", HasSSBS)

clang/lib/CodeGen/CGBuiltin.cpp

@@ -19037,15 +19037,32 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
     ASTContext::GetBuiltinTypeError Error;
     getContext().GetBuiltinType(BuiltinID, Error, &ICEArguments);
     assert(Error == ASTContext::GE_None && "Should not codegen an error");
+    llvm::Type *DataTy = ConvertType(E->getArg(0)->getType());
+    unsigned Size = DataTy->getPrimitiveSizeInBits();
+    llvm::Type *IntTy =
+        llvm::IntegerType::get(Builder.getContext(), std::max(Size, 32u));
+    Function *F = CGM.getIntrinsic(Intrinsic::amdgcn_update_dpp, IntTy);
+    assert(E->getNumArgs() == 5 || E->getNumArgs() == 6);
+    bool InsertOld = E->getNumArgs() == 5;
+    if (InsertOld)
+      Args.push_back(llvm::PoisonValue::get(IntTy));
     for (unsigned I = 0; I != E->getNumArgs(); ++I) {
-      Args.push_back(EmitScalarOrConstFoldImmArg(ICEArguments, I, E));
+      llvm::Value *V = EmitScalarOrConstFoldImmArg(ICEArguments, I, E);
+      if (I <= !InsertOld && Size < 32) {
+        if (!DataTy->isIntegerTy())
+          V = Builder.CreateBitCast(
+              V, llvm::IntegerType::get(Builder.getContext(), Size));
+        V = Builder.CreateZExtOrBitCast(V, IntTy);
+      }
+      llvm::Type *ExpTy =
+          F->getFunctionType()->getFunctionParamType(I + InsertOld);
+      Args.push_back(Builder.CreateTruncOrBitCast(V, ExpTy));
     }
-    assert(Args.size() == 5 || Args.size() == 6);
-    if (Args.size() == 5)
-      Args.insert(Args.begin(), llvm::PoisonValue::get(Args[0]->getType()));
-    Function *F =
-        CGM.getIntrinsic(Intrinsic::amdgcn_update_dpp, Args[0]->getType());
-    return Builder.CreateCall(F, Args);
+    Value *V = Builder.CreateCall(F, Args);
+    if (Size < 32 && !DataTy->isIntegerTy())
+      V = Builder.CreateTrunc(
+          V, llvm::IntegerType::get(Builder.getContext(), Size));
+    return Builder.CreateTruncOrBitCast(V, DataTy);
   }
   case AMDGPU::BI__builtin_amdgcn_permlane16:
   case AMDGPU::BI__builtin_amdgcn_permlanex16:

clang/lib/Format/UnwrappedLineParser.cpp

@@ -2148,7 +2148,7 @@ bool UnwrappedLineParser::tryToParsePropertyAccessor() {
   if (!Style.isCSharp())
     return false;
   // See if it's a property accessor.
-  if (FormatTok->Previous->isNot(tok::identifier))
+  if (!FormatTok->Previous || FormatTok->Previous->isNot(tok::identifier))
     return false;
 
   // See if we are inside a property accessor.

clang/lib/Headers/opencl-c-base.h

@@ -46,6 +46,7 @@
 #define __opencl_c_ext_fp32_global_atomic_min_max 1
 #define __opencl_c_ext_fp32_local_atomic_min_max 1
 #define __opencl_c_ext_image_raw10_raw12 1
+#define __opencl_c_ext_image_unorm_int_2_101010 1
 #define cl_khr_kernel_clock 1
 #define __opencl_c_kernel_clock_scope_device 1
 #define __opencl_c_kernel_clock_scope_work_group 1
@@ -486,6 +487,9 @@ typedef enum memory_order
 #define CLK_UNSIGNED_INT_RAW10_EXT 0x10E3
 #define CLK_UNSIGNED_INT_RAW12_EXT 0x10E4
 #endif // __opencl_c_ext_image_raw10_raw12
+#ifdef __opencl_c_ext_image_unorm_int_2_101010
+#define CLK_UNORM_INT_2_101010_EXT 0x10E5
+#endif // __opencl_c_ext_image_unorm_int_2_101010
 
 // Channel order, numbering must be aligned with cl_channel_order in cl.h
 //

clang/lib/Sema/SemaAMDGPU.cpp

@@ -63,6 +63,44 @@ bool SemaAMDGPU::CheckAMDGCNBuiltinFunctionCall(unsigned BuiltinID,
     OrderIndex = 0;
     ScopeIndex = 1;
     break;
+  case AMDGPU::BI__builtin_amdgcn_mov_dpp: {
+    if (SemaRef.checkArgCountRange(TheCall, 5, 5))
+      return true;
+    Expr *ValArg = TheCall->getArg(0);
+    QualType Ty = ValArg->getType();
+    // TODO: Vectors can also be supported.
+    if (!Ty->isArithmeticType() || Ty->isAnyComplexType()) {
+      SemaRef.Diag(ValArg->getBeginLoc(),
+                   diag::err_typecheck_cond_expect_int_float)
+          << Ty << ValArg->getSourceRange();
+      return true;
+    }
+    return false;
+  }
+  case AMDGPU::BI__builtin_amdgcn_update_dpp: {
+    if (SemaRef.checkArgCountRange(TheCall, 6, 6))
+      return true;
+    Expr *Args[2];
+    QualType ArgTys[2];
+    for (unsigned I = 0; I != 2; ++I) {
+      Args[I] = TheCall->getArg(I);
+      ArgTys[I] = Args[I]->getType();
+      // TODO: Vectors can also be supported.
+      if (!ArgTys[I]->isArithmeticType() || ArgTys[I]->isAnyComplexType()) {
+        SemaRef.Diag(Args[I]->getBeginLoc(),
+                     diag::err_typecheck_cond_expect_int_float)
+            << ArgTys[I] << Args[I]->getSourceRange();
+        return true;
+      }
+    }
+    if (ArgTys[0] != ArgTys[1]) {
+      SemaRef.Diag(Args[1]->getBeginLoc(),
+                   diag::err_typecheck_call_different_arg_types)
+          << ArgTys[0] << ArgTys[1];
+      return true;
+    }
+    return false;
+  }
   default:
     return false;
   }

clang/lib/Sema/SemaOpenMP.cpp

@@ -18335,7 +18335,8 @@ static bool checkOMPArraySectionConstantForReduction(
         return false;
 
       // This is an array subscript which has implicit length 1!
-      ArraySizes.push_back(llvm::APSInt::get(1));
+      llvm::APSInt ConstantOne = llvm::APSInt::get(1);
+      ArraySizes.push_back(ConstantOne);
     } else {
       Expr::EvalResult Result;
       if (!Length->EvaluateAsInt(Result, Context))
@@ -18354,7 +18355,8 @@ static bool checkOMPArraySectionConstantForReduction(
   if (!SingleElement) {
     while (const auto *TempASE = dyn_cast<ArraySubscriptExpr>(Base)) {
       // Has implicit length 1!
-      ArraySizes.push_back(llvm::APSInt::get(1));
+      llvm::APSInt ConstantOne = llvm::APSInt::get(1);
+      ArraySizes.push_back(ConstantOne);
       Base = TempASE->getBase()->IgnoreParenImpCasts();
     }
   }

clang/test/CodeGen/aarch64-cpu-supports-target.c

@@ -17,7 +17,7 @@ int check_all_feature() {
     return 7;
   else if (__builtin_cpu_supports("sve2-bitperm+sve2-sha3+sve2-sm4"))
     return 8;
-  else if (__builtin_cpu_supports("sme+memtag+memtag2+memtag3+sb"))
+  else if (__builtin_cpu_supports("sme+memtag+memtag3+sb"))
     return 9;
   else if (__builtin_cpu_supports("predres+ssbs+ssbs2+bti+ls64+ls64_v"))
     return 10;

clang/test/CodeGen/aarch64-cpu-supports.c

@@ -1,9 +1,10 @@
-// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --check-globals --version 2
+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --check-globals --global-value-regex ".*"
 // RUN: %clang_cc1 -triple aarch64-none-linux-gnu -emit-llvm -o - %s | FileCheck %s
 
+//.
 // CHECK: @__aarch64_cpu_features = external dso_local global { i64 }
-// CHECK-LABEL: define dso_local i32 @main
-// CHECK-SAME: () #[[ATTR0:[0-9]+]] {
+//.
+// CHECK-LABEL: @main(
 // CHECK-NEXT:  entry:
 // CHECK-NEXT:    [[RETVAL:%.*]] = alloca i32, align 4
 // CHECK-NEXT:    store i32 0, ptr [[RETVAL]], align 4
@@ -17,8 +18,8 @@
 // CHECK-NEXT:    br label [[RETURN:%.*]]
 // CHECK:       if.end:
 // CHECK-NEXT:    [[TMP4:%.*]] = load i64, ptr @__aarch64_cpu_features, align 8
-// CHECK-NEXT:    [[TMP5:%.*]] = and i64 [[TMP4]], 9070970929152
-// CHECK-NEXT:    [[TMP6:%.*]] = icmp eq i64 [[TMP5]], 9070970929152
+// CHECK-NEXT:    [[TMP5:%.*]] = and i64 [[TMP4]], 17867063951360
+// CHECK-NEXT:    [[TMP6:%.*]] = icmp eq i64 [[TMP5]], 17867063951360
 // CHECK-NEXT:    [[TMP7:%.*]] = and i1 true, [[TMP6]]
 // CHECK-NEXT:    br i1 [[TMP7]], label [[IF_THEN1:%.*]], label [[IF_END2:%.*]]
 // CHECK:       if.then1:
@@ -60,3 +61,7 @@ int main(void) {
 
   return 0;
 }
+//.
+// CHECK: [[META0:![0-9]+]] = !{i32 1, !"wchar_size", i32 4}
+// CHECK: [[META1:![0-9]+]] = !{!"{{.*}}clang version {{.*}}"}
+//.