Merge from 'main' to 'sycl-web' (58 commits)

  CONFLICT (content): Merge conflict in clang/include/clang/Basic/BuiltinsAMDGPU.def

Author: mdfazlay

clang/include/clang/Basic/BuiltinsAMDGPU.def

@@ -411,5 +411,21 @@ TARGET_BUILTIN(__builtin_amdgcn_cvt_sr_fp8_f32, "ifiiIi", "nc", "fp8-insts")
 //===----------------------------------------------------------------------===//
 BUILTIN(__builtin_amdgcn_implicit_offset, "Ui*5", "nc")
 
+//===----------------------------------------------------------------------===//
+// GFX12+ only builtins.
+//===----------------------------------------------------------------------===//
+
+TARGET_BUILTIN(__builtin_amdgcn_s_barrier_signal, "vIi", "n", "gfx12-insts")
+TARGET_BUILTIN(__builtin_amdgcn_s_barrier_signal_var, "vi", "n", "gfx12-insts")
+TARGET_BUILTIN(__builtin_amdgcn_s_barrier_wait, "vIs", "n", "gfx12-insts")
+TARGET_BUILTIN(__builtin_amdgcn_s_barrier_signal_isfirst, "bIi", "n", "gfx12-insts")
+TARGET_BUILTIN(__builtin_amdgcn_s_barrier_signal_isfirst_var, "bi", "n", "gfx12-insts")
+TARGET_BUILTIN(__builtin_amdgcn_s_barrier_init, "vii", "n", "gfx12-insts")
+TARGET_BUILTIN(__builtin_amdgcn_s_barrier_join, "vi", "n", "gfx12-insts")
+TARGET_BUILTIN(__builtin_amdgcn_s_wakeup_barrier, "vi", "n", "gfx12-insts")
+TARGET_BUILTIN(__builtin_amdgcn_s_barrier_leave, "b", "n", "gfx12-insts")
+TARGET_BUILTIN(__builtin_amdgcn_s_get_barrier_state, "Uii", "n", "gfx12-insts")
+
+
 #undef BUILTIN
 #undef TARGET_BUILTIN

clang/lib/AST/Interp/Interp.cpp

@@ -350,11 +350,6 @@ bool CheckCallable(InterpState &S, CodePtr OpPC, const Function *F) {
   }
 
   if (!F->isConstexpr()) {
-    // Don't emit anything if we're checking for a potential constant
-    // expression. That will happen later when actually executing.
-    if (S.checkingPotentialConstantExpression())
-      return false;
-
     const SourceLocation &Loc = S.Current->getLocation(OpPC);
     if (S.getLangOpts().CPlusPlus11) {
       const FunctionDecl *DiagDecl = F->getDecl();
@@ -371,13 +366,21 @@ bool CheckCallable(InterpState &S, CodePtr OpPC, const Function *F) {
       // FIXME: If DiagDecl is an implicitly-declared special member function
       // or an inheriting constructor, we should be much more explicit about why
       // it's not constexpr.
-      if (CD && CD->isInheritingConstructor())
+      if (CD && CD->isInheritingConstructor()) {
         S.FFDiag(Loc, diag::note_constexpr_invalid_inhctor, 1)
           << CD->getInheritedConstructor().getConstructor()->getParent();
-      else
+        S.Note(DiagDecl->getLocation(), diag::note_declared_at);
+      } else {
+        // Don't emit anything if the function isn't defined and we're checking
+        // for a constant expression. It might be defined at the point we're
+        // actually calling it.
+        if (!DiagDecl->isDefined() && S.checkingPotentialConstantExpression())
+          return false;
+
         S.FFDiag(Loc, diag::note_constexpr_invalid_function, 1)
           << DiagDecl->isConstexpr() << (bool)CD << DiagDecl;
-      S.Note(DiagDecl->getLocation(), diag::note_declared_at);
+        S.Note(DiagDecl->getLocation(), diag::note_declared_at);
+      }
     } else {
       S.FFDiag(Loc, diag::note_invalid_subexpr_in_const_expr);
     }

clang/lib/Headers/CMakeLists.txt

@@ -387,6 +387,8 @@ if(ARM IN_LIST LLVM_TARGETS_TO_BUILD OR AArch64 IN_LIST LLVM_TARGETS_TO_BUILD)
   clang_generate_header(-gen-arm-mve-header arm_mve.td arm_mve.h)
   # Generate arm_cde.h
   clang_generate_header(-gen-arm-cde-header arm_cde.td arm_cde.h)
+  # Generate arm_vector_types.h
+  clang_generate_header(-gen-arm-vector-type arm_neon.td arm_vector_types.h)
 
   # Add headers to target specific lists
   list(APPEND arm_common_generated_files
@@ -403,6 +405,7 @@ if(ARM IN_LIST LLVM_TARGETS_TO_BUILD OR AArch64 IN_LIST LLVM_TARGETS_TO_BUILD)
     "${CMAKE_CURRENT_BINARY_DIR}/arm_sve.h"
     "${CMAKE_CURRENT_BINARY_DIR}/arm_sme_draft_spec_subject_to_change.h"
     "${CMAKE_CURRENT_BINARY_DIR}/arm_bf16.h"
+    "${CMAKE_CURRENT_BINARY_DIR}/arm_vector_types.h"
     )
 endif()
 if(RISCV IN_LIST LLVM_TARGETS_TO_BUILD)

clang/lib/Sema/SemaDecl.cpp

@@ -15913,7 +15913,7 @@ Decl *Sema::ActOnSkippedFunctionBody(Decl *Decl) {
 }
 
 Decl *Sema::ActOnFinishFunctionBody(Decl *D, Stmt *BodyArg) {
-  return ActOnFinishFunctionBody(D, BodyArg, false);
+  return ActOnFinishFunctionBody(D, BodyArg, /*IsInstantiation=*/false);
 }
 
 /// RAII object that pops an ExpressionEvaluationContext when exiting a function

clang/lib/Sema/SemaType.cpp

@@ -8447,12 +8447,25 @@ static void HandleNeonVectorTypeAttr(QualType &CurType, const ParsedAttr &Attr,
   // not to need a separate attribute)
   if (!(S.Context.getTargetInfo().hasFeature("neon") ||
         S.Context.getTargetInfo().hasFeature("mve") ||
-        IsTargetCUDAAndHostARM)) {
+        S.Context.getTargetInfo().hasFeature("sve") ||
+        S.Context.getTargetInfo().hasFeature("sme") ||
+        IsTargetCUDAAndHostARM) &&
+      VecKind == VectorKind::Neon) {
+    S.Diag(Attr.getLoc(), diag::err_attribute_unsupported)
+        << Attr << "'neon', 'mve', 'sve' or 'sme'";
+    Attr.setInvalid();
+    return;
+  }
+  if (!(S.Context.getTargetInfo().hasFeature("neon") ||
+        S.Context.getTargetInfo().hasFeature("mve") ||
+        IsTargetCUDAAndHostARM) &&
+      VecKind == VectorKind::NeonPoly) {
     S.Diag(Attr.getLoc(), diag::err_attribute_unsupported)
         << Attr << "'neon' or 'mve'";
     Attr.setInvalid();
     return;
   }
+
   // Check the attribute arguments.
   if (Attr.getNumArgs() != 1) {
     S.Diag(Attr.getLoc(), diag::err_attribute_wrong_number_arguments)

clang/test/AST/Interp/functions.cpp

@@ -267,6 +267,17 @@ namespace InvalidCall {
                    // ref-error {{must be initialized by a constant expression}} \
                    // ref-note {{in call to 'SS()'}}
 
+
+  /// This should not emit a diagnostic.
+  constexpr int f();
+  constexpr int a() {
+    return f();
+  }
+  constexpr int f() {
+    return 5;
+  }
+  static_assert(a() == 5, "");
+
 }
 
 namespace CallWithArgs {

clang/test/CodeGen/arm-vector_type-params-returns.c

@@ -0,0 +1,136 @@
+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 3
+
+// RUN: %clang_cc1 -DSVE_HEADER -triple aarch64 -target-feature +sve -emit-llvm -O2 -o - %s | opt -S -passes=mem2reg,sroa | FileCheck %s
+// RUN: %clang_cc1 -DSVE_HEADER -triple aarch64-none-linux-gnu -target-feature +sve2p1 -S -disable-O0-optnone -Werror -Wall -o - /dev/null %s
+
+// RUN: %clang_cc1 -DNEON_HEADER -triple aarch64 -target-feature +sve -emit-llvm -O2 -o - %s | opt -S -passes=mem2reg,sroa | FileCheck %s
+// RUN: %clang_cc1 -DNEON_HEADER -triple aarch64-none-linux-gnu -target-feature +sve2p1 -S -disable-O0-optnone -Werror -Wall -o - /dev/null %s
+
+// RUN: %clang_cc1 -DSVE_HEADER -DNEON_HEADER -triple aarch64 -target-feature +sve -emit-llvm -O2 -o - %s | opt -S -passes=mem2reg,sroa | FileCheck %s
+// RUN: %clang_cc1 -DSVE_HEADER -DNEON_HEADER -triple aarch64-none-linux-gnu -target-feature +sve2p1 -S -disable-O0-optnone -Werror -Wall -o - /dev/null %s
+
+// RUN: %clang_cc1 -DNEON_HEADER -DSVE_HEADER2  -triple aarch64 -target-feature +sve -emit-llvm -O2 -o - %s | opt -S -passes=mem2reg,sroa | FileCheck %s
+// RUN: %clang_cc1 -DNEON_HEADER -DSVE_HEADER2 -triple aarch64-none-linux-gnu -target-feature +sve2p1 -S -disable-O0-optnone -Werror -Wall -o - /dev/null %s
+
+// REQUIRES: aarch64-registered-target
+
+#ifdef SVE_HEADER
+  #include <arm_sve.h>
+#endif
+
+#ifdef NEON_HEADER
+  #include <arm_neon.h>
+#endif
+
+#ifdef SVE_HEADER_2
+  #include <arm_sve.h>
+#endif
+
+// function return types
+// CHECK-LABEL: define dso_local <8 x half> @test_ret_v8f16(
+// CHECK-SAME: <8 x half> noundef returned [[V:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] {
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    ret <8 x half> [[V]]
+//
+float16x8_t test_ret_v8f16(float16x8_t v) {
+  return v;
+}
+
+// CHECK-LABEL: define dso_local <4 x float> @test_ret_v4f32(
+// CHECK-SAME: <4 x float> noundef returned [[V:%.*]]) local_unnamed_addr #[[ATTR0]] {
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    ret <4 x float> [[V]]
+//
+float32x4_t test_ret_v4f32(float32x4_t v) {
+  return v;
+}
+
+// CHECK-LABEL: define dso_local <2 x double> @test_ret_v2f64(
+// CHECK-SAME: <2 x double> noundef returned [[V:%.*]]) local_unnamed_addr #[[ATTR0]] {
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    ret <2 x double> [[V]]
+//
+float64x2_t test_ret_v2f64(float64x2_t v) {
+  return v;
+}
+
+// CHECK-LABEL: define dso_local <8 x bfloat> @test_ret_v8bf16(
+// CHECK-SAME: <8 x bfloat> noundef returned [[V:%.*]]) local_unnamed_addr #[[ATTR0]] {
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    ret <8 x bfloat> [[V]]
+//
+bfloat16x8_t test_ret_v8bf16(bfloat16x8_t v) {
+  return v;
+}
+
+// CHECK-LABEL: define dso_local <16 x i8> @test_ret_v16s8(
+// CHECK-SAME: <16 x i8> noundef returned [[V:%.*]]) local_unnamed_addr #[[ATTR0]] {
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    ret <16 x i8> [[V]]
+//
+int8x16_t test_ret_v16s8(int8x16_t v) {
+  return v;
+}
+
+// CHECK-LABEL: define dso_local <8 x i16> @test_ret_v8s16(
+// CHECK-SAME: <8 x i16> noundef returned [[V:%.*]]) local_unnamed_addr #[[ATTR0]] {
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    ret <8 x i16> [[V]]
+//
+int16x8_t test_ret_v8s16(int16x8_t v) {
+  return v;
+}
+
+// CHECK-LABEL: define dso_local <4 x i32> @test_ret_v32s4(
+// CHECK-SAME: <4 x i32> noundef returned [[V:%.*]]) local_unnamed_addr #[[ATTR0]] {
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    ret <4 x i32> [[V]]
+//
+int32x4_t test_ret_v32s4(int32x4_t v) {
+  return v;
+}
+
+// CHECK-LABEL: define dso_local <2 x i64> @test_ret_v64s2(
+// CHECK-SAME: <2 x i64> noundef returned [[V:%.*]]) local_unnamed_addr #[[ATTR0]] {
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    ret <2 x i64> [[V]]
+//
+int64x2_t test_ret_v64s2(int64x2_t v) {
+  return v;
+}
+
+// CHECK-LABEL: define dso_local <16 x i8> @test_ret_v16u8(
+// CHECK-SAME: <16 x i8> noundef returned [[V:%.*]]) local_unnamed_addr #[[ATTR0]] {
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    ret <16 x i8> [[V]]
+//
+uint8x16_t test_ret_v16u8(uint8x16_t v) {
+  return v;
+}
+
+// CHECK-LABEL: define dso_local <8 x i16> @test_ret_v8u16(
+// CHECK-SAME: <8 x i16> noundef returned [[V:%.*]]) local_unnamed_addr #[[ATTR0]] {
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    ret <8 x i16> [[V]]
+//
+uint16x8_t test_ret_v8u16(uint16x8_t v) {
+  return v;
+}
+
+// CHECK-LABEL: define dso_local <4 x i32> @test_ret_v32u4(
+// CHECK-SAME: <4 x i32> noundef returned [[V:%.*]]) local_unnamed_addr #[[ATTR0]] {
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    ret <4 x i32> [[V]]
+//
+uint32x4_t test_ret_v32u4(uint32x4_t v) {
+  return v;
+}
+
+// CHECK-LABEL: define dso_local <2 x i64> @test_ret_v64u2(
+// CHECK-SAME: <2 x i64> noundef returned [[V:%.*]]) local_unnamed_addr #[[ATTR0]] {
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    ret <2 x i64> [[V]]
+//
+uint64x2_t test_ret_v64u2(uint64x2_t v) {
+  return v;
+}

clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12-err.cl

@@ -0,0 +1,24 @@
+// REQUIRES: amdgpu-registered-target
+
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1200 -verify -S -emit-llvm -o - %s
+
+kernel void builtins_amdgcn_s_barrier_signal_err(global int* in, global int* out, int barrier) {
+
+  __builtin_amdgcn_s_barrier_signal(barrier); // expected-error {{'__builtin_amdgcn_s_barrier_signal' must be a constant integer}}
+  __builtin_amdgcn_s_barrier_wait(-1);
+  *out = *in;
+}
+
+kernel void builtins_amdgcn_s_barrier_wait_err(global int* in, global int* out, int barrier) {
+
+  __builtin_amdgcn_s_barrier_signal(-1);
+  __builtin_amdgcn_s_barrier_wait(barrier); // expected-error {{'__builtin_amdgcn_s_barrier_wait' must be a constant integer}}
+  *out = *in;
+}
+
+kernel void builtins_amdgcn_s_barrier_signal_isfirst_err(global int* in, global int* out, int barrier) {
+
+  __builtin_amdgcn_s_barrier_signal_isfirst(barrier); // expected-error {{'__builtin_amdgcn_s_barrier_signal_isfirst' must be a constant integer}}
+  __builtin_amdgcn_s_barrier_wait(-1);
+  *out = *in;
+}