Merged master:ab6779bbd8f into amd-gfx:3e53220dbf6

Local branch amd-gfx 3e53220 Merged master:116e38fd8b8 into amd-gfx:ac9866a082e
Remote branch master ab6779b [Statepoint] Remove last of old ImmutableStatepoint code

clang/include/clang/Sema/Sema.h

@@ -11724,6 +11724,10 @@ class Sema final {
   void maybeAddCUDAHostDeviceAttrs(FunctionDecl *FD,
                                    const LookupResult &Previous);
 
+  /// May add implicit CUDAConstantAttr attribute to VD, depending on VD
+  /// and current compilation settings.
+  void MaybeAddCUDAConstantAttr(VarDecl *VD);
+
 public:
   /// Check whether we're allowed to call Callee from the current context.
   ///

clang/lib/Sema/SemaCUDA.cpp

@@ -513,9 +513,14 @@ void Sema::checkAllowedCUDAInitializer(VarDecl *VD) {
     // constructor according to CUDA rules. This deviates from NVCC,
     // but allows us to handle things like constexpr constructors.
     if (!AllowedInit &&
-        (VD->hasAttr<CUDADeviceAttr>() || VD->hasAttr<CUDAConstantAttr>()))
-      AllowedInit = VD->getInit()->isConstantInitializer(
-          Context, VD->getType()->isReferenceType());
+        (VD->hasAttr<CUDADeviceAttr>() || VD->hasAttr<CUDAConstantAttr>())) {
+      auto *Init = VD->getInit();
+      AllowedInit =
+          ((VD->getType()->isDependentType() || Init->isValueDependent()) &&
+           VD->isConstexpr()) ||
+          Init->isConstantInitializer(Context,
+                                      VD->getType()->isReferenceType());
+    }
 
     // Also make sure that destructor, if there is one, is empty.
     if (AllowedInit)
@@ -612,6 +617,13 @@ void Sema::maybeAddCUDAHostDeviceAttrs(FunctionDecl *NewD,
   NewD->addAttr(CUDADeviceAttr::CreateImplicit(Context));
 }
 
+void Sema::MaybeAddCUDAConstantAttr(VarDecl *VD) {
+  if (getLangOpts().CUDAIsDevice && VD->isConstexpr() &&
+      (VD->isFileVarDecl() || VD->isStaticDataMember())) {
+    VD->addAttr(CUDAConstantAttr::CreateImplicit(getASTContext()));
+  }
+}
+
 Sema::DeviceDiagBuilder Sema::CUDADiagIfDeviceCode(SourceLocation Loc,
                                                    unsigned DiagID) {
   assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");

clang/lib/Sema/SemaDecl.cpp

@@ -7100,6 +7100,7 @@ NamedDecl *Sema::ActOnVariableDeclarator(
 
   case CSK_constexpr:
     NewVD->setConstexpr(true);
+    MaybeAddCUDAConstantAttr(NewVD);
     // C++1z [dcl.spec.constexpr]p1:
     //   A static data member declared with the constexpr specifier is
     //   implicitly an inline variable.

clang/lib/Sema/SemaTemplateInstantiateDecl.cpp

@@ -4841,6 +4841,7 @@ void Sema::BuildVariableInstantiation(
   NewVar->setCXXForRangeDecl(OldVar->isCXXForRangeDecl());
   NewVar->setObjCForDecl(OldVar->isObjCForDecl());
   NewVar->setConstexpr(OldVar->isConstexpr());
+  MaybeAddCUDAConstantAttr(NewVar);
   NewVar->setInitCapture(OldVar->isInitCapture());
   NewVar->setPreviousDeclInSameBlockScope(
       OldVar->isPreviousDeclInSameBlockScope());

clang/test/CodeGenCUDA/constexpr-variables.cu

@@ -0,0 +1,43 @@
+// RUN: %clang_cc1 -std=c++14 %s -emit-llvm -o - -triple nvptx \
+// RUN:   -fcuda-is-device | FileCheck --check-prefixes=CXX14 %s
+// RUN: %clang_cc1 -std=c++17 %s -emit-llvm -o - -triple nvptx \
+// RUN:   -fcuda-is-device | FileCheck --check-prefixes=CXX17 %s
+
+#include "Inputs/cuda.h"
+
+// COM: @_ZL1a = internal {{.*}}constant i32 7
+constexpr int a = 7;
+__constant__ const int &use_a = a;
+
+namespace B {
+ // COM: @_ZN1BL1bE = internal {{.*}}constant i32 9
+  constexpr int b = 9;
+}
+__constant__ const int &use_B_b = B::b;
+
+struct Q {
+  // CXX14: @_ZN1Q2k2E = {{.*}}externally_initialized constant i32 6
+  // CXX17: @_ZN1Q2k2E = internal {{.*}}constant i32 6
+  // CXX14: @_ZN1Q2k1E = available_externally {{.*}}constant i32 5
+  // CXX17: @_ZN1Q2k1E = linkonce_odr {{.*}}constant i32 5
+  static constexpr int k1 = 5;
+  static constexpr int k2 = 6;
+};
+constexpr int Q::k2;
+
+__constant__ const int &use_Q_k1 = Q::k1;
+__constant__ const int &use_Q_k2 = Q::k2;
+
+template<typename T> struct X {
+  // CXX14: @_ZN1XIiE1aE = available_externally {{.*}}constant i32 123
+  // CXX17: @_ZN1XIiE1aE = linkonce_odr {{.*}}constant i32 123
+  static constexpr int a = 123;
+};
+__constant__ const int &use_X_a = X<int>::a;
+
+template <typename T, T a, T b> struct A {
+  // CXX14: @_ZN1AIiLi1ELi2EE1xE = available_externally {{.*}}constant i32 2
+  // CXX17: @_ZN1AIiLi1ELi2EE1xE = linkonce_odr {{.*}}constant i32 2
+  constexpr static T x = a * b;
+};
+__constant__ const int &y = A<int, 1, 2>::x;

clang/test/SemaCUDA/constexpr-variables.cu

@@ -0,0 +1,80 @@
+// RUN: %clang_cc1 -std=c++14 %s -emit-llvm -o - -triple nvptx64-nvidia-cuda \
+// RUN:   -fcuda-is-device -verify -fsyntax-only
+// RUN: %clang_cc1 -std=c++17 %s -emit-llvm -o - -triple nvptx64-nvidia-cuda \
+// RUN:   -fcuda-is-device -verify -fsyntax-only
+// RUN: %clang_cc1 -std=c++14 %s -emit-llvm -o - \
+// RUN:   -triple x86_64-unknown-linux-gnu -verify -fsyntax-only
+// RUN: %clang_cc1 -std=c++17 %s -emit-llvm -o - \
+// RUN:   -triple x86_64-unknown-linux-gnu -verify -fsyntax-only
+#include "Inputs/cuda.h"
+
+template<typename T>
+__host__ __device__ void foo(const T **a) {
+  // expected-note@-1 {{declared here}}
+  static const T b = sizeof(a);
+  static constexpr T c = sizeof(a);
+  const T d = sizeof(a);
+  constexpr T e = sizeof(a);
+  constexpr T f = **a;
+  // expected-error@-1 {{constexpr variable 'f' must be initialized by a constant expression}}
+  // expected-note@-2 {{read of non-constexpr variable 'a' is not allowed in a constant expression}}
+  a[0] = &b;
+  a[1] = &c;
+  a[2] = &d;
+  a[3] = &e;
+}
+
+__device__ void device_fun(const int **a) {
+  // expected-note@-1 {{declared here}}
+  constexpr int b = sizeof(a);
+  static constexpr int c = sizeof(a);
+  constexpr int d = **a;
+  // expected-error@-1 {{constexpr variable 'd' must be initialized by a constant expression}}
+  // expected-note@-2 {{read of non-constexpr variable 'a' is not allowed in a constant expression}}
+  a[0] = &b;
+  a[1] = &c;
+  foo(a);
+  // expected-note@-1 {{in instantiation of function template specialization 'foo<int>' requested here}}
+}
+
+void host_fun(const int **a) {
+  // expected-note@-1 {{declared here}}
+  constexpr int b = sizeof(a);
+  static constexpr int c = sizeof(a);
+  constexpr int d = **a;
+  // expected-error@-1 {{constexpr variable 'd' must be initialized by a constant expression}}
+  // expected-note@-2 {{read of non-constexpr variable 'a' is not allowed in a constant expression}}
+  a[0] = &b;
+  a[1] = &c;
+  foo(a);
+}
+
+__host__ __device__ void host_device_fun(const int **a) {
+  // expected-note@-1 {{declared here}}
+  constexpr int b = sizeof(a);
+  static constexpr int c = sizeof(a);
+  constexpr int d = **a;
+  // expected-error@-1 {{constexpr variable 'd' must be initialized by a constant expression}}
+  // expected-note@-2 {{read of non-constexpr variable 'a' is not allowed in a constant expression}}
+  a[0] = &b;
+  a[1] = &c;
+  foo(a);
+}
+
+template <class T>
+struct A {
+  explicit A() = default;
+};
+template <class T>
+constexpr A<T> a{};
+
+struct B {
+  static constexpr bool value = true;
+};
+
+template<typename T>
+struct C {
+  static constexpr bool value = T::value;
+};
+
+__constant__ const bool &x = C<B>::value;

compiler-rt/lib/sanitizer_common/sanitizer_mac.cpp

@@ -27,7 +27,6 @@
 #include "sanitizer_flags.h"
 #include "sanitizer_internal_defs.h"
 #include "sanitizer_libc.h"
-#include "sanitizer_placement_new.h"
 #include "sanitizer_platform_limits_posix.h"
 #include "sanitizer_procmaps.h"
 #include "sanitizer_ptrauth.h"
@@ -630,19 +629,24 @@ MacosVersion GetMacosAlignedVersion() {
   return *reinterpret_cast<MacosVersion *>(&result);
 }
 
+void ParseVersion(const char *vers, u16 *major, u16 *minor) {
+  // Format: <major>.<minor>.<patch>\0
+  CHECK_GE(internal_strlen(vers), 5);
+  const char *p = vers;
+  *major = internal_simple_strtoll(p, &p, /*base=*/10);
+  CHECK_EQ(*p, '.');
+  p += 1;
+  *minor = internal_simple_strtoll(p, &p, /*base=*/10);
+}
+
 DarwinKernelVersion GetDarwinKernelVersion() {
   char buf[100];
   size_t len = sizeof(buf);
   int res = internal_sysctlbyname("kern.osrelease", buf, &len, nullptr, 0);
   CHECK_EQ(res, 0);
 
-  // Format: <major>.<minor>.<patch>\0
-  CHECK_GE(len, 6);
-  const char *p = buf;
-  u16 major = internal_simple_strtoll(p, &p, /*base=*/10);
-  CHECK_EQ(*p, '.');
-  p += 1;
-  u16 minor = internal_simple_strtoll(p, &p, /*base=*/10);
+  u16 major, minor;
+  ParseVersion(buf, &major, &minor);
 
   return DarwinKernelVersion(major, minor);
 }
@@ -840,9 +844,9 @@ bool DyldNeedsEnvVariable() {
   if (!&dyldVersionNumber) return true;
   // If running on OS X 10.11+ or iOS 9.0+, dyld will interpose even if
   // DYLD_INSERT_LIBRARIES is not set. However, checking OS version via
-  // GetMacosAlignedVersion() doesn't work for the simulator. Let's instead check
-  // `dyldVersionNumber`, which is exported by dyld, against a known version
-  // number from the first OS release where this appeared.
+  // GetMacosAlignedVersion() doesn't work for the simulator. Let's instead
+  // check `dyldVersionNumber`, which is exported by dyld, against a known
+  // version number from the first OS release where this appeared.
   return dyldVersionNumber < kMinDyldVersionWithAutoInterposition;
 }
 

compiler-rt/lib/sanitizer_common/tests/CMakeLists.txt

@@ -21,6 +21,7 @@ set(SANITIZER_UNITTESTS
   sanitizer_libc_test.cpp
   sanitizer_linux_test.cpp
   sanitizer_list_test.cpp
+  sanitizer_mac_test.cpp
   sanitizer_mutex_test.cpp
   sanitizer_nolibc_test.cpp
   sanitizer_posix_test.cpp

compiler-rt/lib/sanitizer_common/tests/sanitizer_mac_test.cpp

@@ -0,0 +1,58 @@
+//===-- sanitizer_mac_test.cpp --------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// Tests for sanitizer_mac.{h,cpp}
+//
+//===----------------------------------------------------------------------===//
+
+#include "sanitizer_common/sanitizer_platform.h"
+#if SANITIZER_MAC
+
+#include "sanitizer_common/sanitizer_mac.h"
+
+#include "gtest/gtest.h"
+
+#include <sys/sysctl.h>  // sysctlbyname
+#include <mach/kern_return.h>  // KERN_SUCCESS
+
+namespace __sanitizer {
+
+TEST(SanitizerMac, GetMacosAlignedVersion) {
+  MacosVersion vers = GetMacosAlignedVersion();
+  EXPECT_EQ(vers.major, 10);
+  EXPECT_EQ(vers.minor, GetDarwinKernelVersion().major - 4);
+}
+
+void ParseVersion(const char *vers, u16 *major, u16 *minor);
+
+TEST(SanitizerMac, ParseVersion) {
+  u16 major, minor;
+  ParseVersion("11.22.33", &major, &minor);
+  EXPECT_EQ(major, 11);
+  EXPECT_EQ(minor, 22);
+}
+
+TEST(SanitizerMac, GetDarwinKernelVersion) {
+  DarwinKernelVersion vers = GetDarwinKernelVersion();
+  std::ostringstream oss;
+  oss << vers.major << '.' << vers.minor;
+  std::string actual = oss.str();
+
+  char buf[100];
+  size_t len = sizeof(buf);
+  int res = sysctlbyname("kern.osrelease", buf, &len, nullptr, 0);
+  ASSERT_EQ(res, KERN_SUCCESS);
+  std::string expected(buf);
+
+  // Prefix match
+  ASSERT_TRUE(expected.compare(0, actual.size(), actual) == 0);
+}
+
+}  // namespace __sanitizer
+
+#endif  // SANITIZER_MAC

llvm/include/llvm/CodeGen/GlobalISel/CallLowering.h

@@ -61,7 +61,8 @@ class CallLowering {
       if (!Regs.empty() && Flags.empty())
         this->Flags.push_back(ISD::ArgFlagsTy());
       // FIXME: We should have just one way of saying "no register".
-      assert((Ty->isVoidTy() == (Regs.empty() || Regs[0] == 0)) &&
+      assert(((Ty->isVoidTy() || Ty->isEmptyTy()) ==
+              (Regs.empty() || Regs[0] == 0)) &&
              "only void types should have no register");
     }
 

llvm/include/llvm/CodeGen/GlobalISel/IRTranslator.h

@@ -236,6 +236,9 @@ class IRTranslator : public MachineFunctionPass {
   bool translateSimpleIntrinsic(const CallInst &CI, Intrinsic::ID ID,
                                 MachineIRBuilder &MIRBuilder);
 
+  bool translateConstrainedFPIntrinsic(const ConstrainedFPIntrinsic &FPI,
+                                       MachineIRBuilder &MIRBuilder);
+
   bool translateKnownIntrinsic(const CallInst &CI, Intrinsic::ID ID,
                                MachineIRBuilder &MIRBuilder);