Merge branch 'main' into cfgsimplify

Author: luciechoi

clang/lib/CodeGen/CGPointerAuth.cpp

@@ -440,9 +440,10 @@ CodeGenModule::getConstantSignedPointer(llvm::Constant *Pointer, unsigned Key,
     IntegerDiscriminator = llvm::ConstantInt::get(Int64Ty, 0);
   }
 
-  return llvm::ConstantPtrAuth::get(Pointer,
-                                    llvm::ConstantInt::get(Int32Ty, Key),
-                                    IntegerDiscriminator, AddressDiscriminator);
+  return llvm::ConstantPtrAuth::get(
+      Pointer, llvm::ConstantInt::get(Int32Ty, Key), IntegerDiscriminator,
+      AddressDiscriminator,
+      /*DeactivationSymbol=*/llvm::Constant::getNullValue(DefaultPtrTy));
 }
 
 /// Does a given PointerAuthScheme require us to sign a value

clang/lib/Driver/ToolChains/Linux.cpp

@@ -922,7 +922,7 @@ SanitizerMask Linux::getSupportedSanitizers() const {
   if (IsX86_64 || IsMIPS64 || IsAArch64 || IsPowerPC64 || IsSystemZ ||
       IsLoongArch64 || IsRISCV64)
     Res |= SanitizerKind::Thread;
-  if (IsX86_64 || IsAArch64 || IsSystemZ)
+  if (IsX86_64 || IsAArch64)
     Res |= SanitizerKind::Type;
   if (IsX86_64 || IsSystemZ || IsPowerPC64)
     Res |= SanitizerKind::KernelMemory;

clang/test/Driver/hip-spirv-backend-opt.c

@@ -4,58 +4,58 @@
 
 // RUN: %clang --offload-new-driver --target=x86_64-unknown-linux-gnu --offload-arch=amdgcnspirv \
 // RUN:         -nogpuinc -nogpulib -### -x hip %s -save-temps  \
-// RUN:         -use-spirv-backend --offload-device-only -S \
+// RUN:         -use-spirv-backend --offload-device-only -S -no-canonical-prefixes \
 // RUN: 2>&1 | FileCheck %s --check-prefixes=CHECK-SPIRV-TRANSLATOR,CHECK-SPIRV-BACKEND-TEXTUAL
 
 // RUN: %clang --offload-new-driver --target=x86_64-unknown-linux-gnu --offload-arch=amdgcnspirv \
 // RUN:         -nogpuinc -nogpulib -### -x hip %s -save-temps  \
-// RUN:         -use-spirv-backend --offload-device-only \
+// RUN:         -use-spirv-backend --offload-device-only -no-canonical-prefixes \
 // RUN: 2>&1 | FileCheck %s --check-prefixes=CHECK-SPIRV-TRANSLATOR,CHECK-SPIRV-BACKEND-BINARY
 
 // The new driver's behavior is to emit LLVM IR for --offload-device-only and -fgpu-rdc (independently of SPIR-V).
 // RUN: %clang --offload-new-driver --target=x86_64-unknown-linux-gnu --offload-arch=amdgcnspirv \
 // RUN:         -### -nogpuinc -nogpulib -x hip %s -save-temps \
-// RUN:         -use-spirv-backend --offload-device-only -S -fgpu-rdc \
+// RUN:         -use-spirv-backend --offload-device-only -S -fgpu-rdc -no-canonical-prefixes \
 // RUN: 2>&1 | FileCheck %s --check-prefixes=CHECK-SPIRV-TRANSLATOR,CHECK-SPIRV-BACKEND-LL,CHECK-FGPU-RDC
 
 // The new driver's behavior is to emit LLVM IR for --offload-device-only and -fgpu-rdc (independently of SPIR-V).
 // RUN: %clang --offload-new-driver --target=x86_64-unknown-linux-gnu --offload-arch=amdgcnspirv \
 // RUN:         -nogpuinc -nogpulib -### -x hip %s -save-temps  \
-// RUN:         -use-spirv-backend --offload-device-only -fgpu-rdc \
+// RUN:         -use-spirv-backend --offload-device-only -fgpu-rdc -no-canonical-prefixes \
 // RUN: 2>&1 | FileCheck %s --check-prefixes=CHECK-SPIRV-TRANSLATOR,CHECK-SPIRV-BACKEND-BC,CHECK-FGPU-RDC
 
 // --offload-device-only is always unset --- testing interactions with -S and -fgpu-rdc
 
 // RUN: %clang --offload-new-driver --target=x86_64-unknown-linux-gnu --offload-arch=amdgcnspirv \
 // RUN:         -nogpuinc -nogpulib -### -x hip %s -save-temps  \
-// RUN:         -use-spirv-backend -S -fgpu-rdc \
+// RUN:         -use-spirv-backend -S -fgpu-rdc -no-canonical-prefixes \
 // RUN: 2>&1 | FileCheck %s --check-prefixes=CHECK-SPIRV-TRANSLATOR,CHECK-SPIRV-BACKEND-BC,CHECK-FGPU-RDC
 
 // RUN: %clang --offload-new-driver --target=x86_64-unknown-linux-gnu --offload-arch=amdgcnspirv \
 // RUN:         -nogpuinc -nogpulib -### -x hip %s -save-temps  \
-// RUN:         -use-spirv-backend -S \
+// RUN:         -use-spirv-backend -S -no-canonical-prefixes \
 // RUN: 2>&1 | FileCheck %s --check-prefixes=CHECK-SPIRV-TRANSLATOR,CHECK-SPIRV-BACKEND-BC
 
 // RUN: %clang --offload-new-driver --target=x86_64-unknown-linux-gnu --offload-arch=amdgcnspirv \
 // RUN:         -nogpuinc -nogpulib -### -x hip %s -save-temps  \
-// RUN:         -use-spirv-backend -fgpu-rdc \
+// RUN:         -use-spirv-backend -fgpu-rdc -no-canonical-prefixes \
 // RUN: 2>&1 | FileCheck %s --check-prefixes=CHECK-SPIRV-TRANSLATOR,CHECK-SPIRV-BACKEND-BC,CHECK-CLANG-LINKER-WRAPPER
 
 // RUN: %clang --offload-new-driver --target=x86_64-unknown-linux-gnu --offload-arch=amdgcnspirv \
 // RUN:         -nogpuinc -nogpulib -### -x hip %s -save-temps  \
-// RUN:         -use-spirv-backend \
+// RUN:         -use-spirv-backend -no-canonical-prefixes \
 // RUN: 2>&1 | FileCheck %s --check-prefixes=CHECK-SPIRV-TRANSLATOR,CHECK-SPIRV-BACKEND-BC,CHECK-CLANG-LINKER-WRAPPER
 
 // RUN: %clang --no-offload-new-driver --target=x86_64-unknown-linux-gnu --offload-arch=amdgcnspirv \
 // RUN:         -nogpuinc -nogpulib -### -x hip %s -save-temps  \
-// RUN:         -use-spirv-backend \
+// RUN:         -use-spirv-backend -no-canonical-prefixes \
 // RUN: 2>&1 | FileCheck %s --check-prefixes=CHECK-SPIRV-TRANSLATOR,CHECK-SPIRV-BACKEND-BC,CHECK-SPIRV-BACKEND-BINARY-EQ-TRIPLE
 
 // CHECK-SPIRV-TRANSLATOR-NOT: "{{.*llvm-spirv.*}}"
-// CHECK-SPIRV-BACKEND-TEXTUAL: "{{.*}}clang{{.*}}" "-cc1" "-triple" "spirv64-amd-amdhsa" {{.*}} "-S"
-// CHECK-SPIRV-BACKEND-BINARY: "{{.*}}clang{{.*}}" "-cc1" "-triple" "spirv64-amd-amdhsa" {{.*}} "-emit-obj"
-// CHECK-SPIRV-BACKEND-BC: "{{.*}}clang{{.*}}" "-cc1" "-triple" "spirv64-amd-amdhsa" {{.*}} "-emit-llvm-bc"
-// CHECK-SPIRV-BACKEND-LL: "{{.*}}clang{{.*}}" "-cc1" "-triple" "spirv64-amd-amdhsa" {{.*}} "-emit-llvm"
-// CHECK-SPIRV-BACKEND-BINARY-EQ-TRIPLE: "{{.*}}clang{{.*}}" "-cc1" {{.*}}"-triple=spirv64-amd-amdhsa" {{.*}}"-emit-obj"
+// CHECK-SPIRV-BACKEND-TEXTUAL: "{{.*clang(\.exe)?}}" "-cc1" "-triple" "spirv64-amd-amdhsa" {{.*}} "-S"
+// CHECK-SPIRV-BACKEND-BINARY: "{{.*clang(\.exe)?}}" "-cc1" "-triple" "spirv64-amd-amdhsa" {{.*}} "-emit-obj"
+// CHECK-SPIRV-BACKEND-BC: "{{.*clang(\.exe)?}}" "-cc1" "-triple" "spirv64-amd-amdhsa" {{.*}} "-emit-llvm-bc"
+// CHECK-SPIRV-BACKEND-LL: "{{.*clang(\.exe)?}}" "-cc1" "-triple" "spirv64-amd-amdhsa" {{.*}} "-emit-llvm"
+// CHECK-SPIRV-BACKEND-BINARY-EQ-TRIPLE: "{{.*clang(\.exe)?}}" "-cc1" {{.*}}"-triple=spirv64-amd-amdhsa" {{.*}}"-emit-obj"
 // CHECK-FGPU-RDC-SAME: {{.*}} "-fgpu-rdc"
 // CHECK-CLANG-LINKER-WRAPPER: "{{.*}}clang-linker-wrapper" "--should-extract=amdgcnspirv" {{.*}} "--device-compiler=spirv64-amd-amdhsa=-use-spirv-backend"

compiler-rt/cmake/Modules/AllSupportedArchDefs.cmake

@@ -89,7 +89,7 @@ else()
   set(ALL_TSAN_SUPPORTED_ARCH ${X86_64} ${MIPS64} ${ARM64} ${PPC64} ${S390X}
       ${LOONGARCH64} ${RISCV64})
 endif()
-set(ALL_TYSAN_SUPPORTED_ARCH ${X86_64} ${ARM64} ${S390X})
+set(ALL_TYSAN_SUPPORTED_ARCH ${X86_64} ${ARM64})
 set(ALL_UBSAN_SUPPORTED_ARCH ${X86} ${X86_64} ${ARM32} ${ARM64} ${RISCV64}
     ${MIPS32} ${MIPS64} ${PPC64} ${S390X} ${SPARC} ${SPARCV9} ${HEXAGON}
     ${LOONGARCH64})

compiler-rt/lib/tysan/tysan_platform.h

@@ -45,12 +45,6 @@ struct Mapping48 {
   static const uptr kPtrShift = 3;
 };
 #define TYSAN_RUNTIME_VMA 1
-#elif defined(__s390x__)
-struct Mapping {
-  static const uptr kShadowAddr = 0x080000000000ULL;
-  static const uptr kAppAddr = 0x460000000000ULL;
-  static const uptr kAppMemMsk = ~0xC00000000000ULL;
-};
 #else
 #error "TySan not supported for this platform!"
 #endif

flang-rt/lib/cuda/allocator.cpp

@@ -19,8 +19,6 @@
 #include "flang/Runtime/CUDA/common.h"
 #include "flang/Support/Fortran.h"
 
-#include "cuda_runtime.h"
-
 namespace Fortran::runtime::cuda {
 
 struct DeviceAllocation {
@@ -133,6 +131,15 @@ void RTDEF(CUFRegisterAllocator)() {
   allocatorRegistry.Register(
       kUnifiedAllocatorPos, {&CUFAllocUnified, CUFFreeUnified});
 }
+
+cudaStream_t RTDECL(CUFAssociatedGetStream)(void *p) {
+  int pos = findAllocation(p);
+  if (pos >= 0) {
+    cudaStream_t stream = deviceAllocations[pos].stream;
+    return stream;
+  }
+  return nullptr;
+}
 }
 
 void *CUFAllocPinned(

flang-rt/unittests/Runtime/CUDA/Allocatable.cpp

@@ -121,3 +121,54 @@ TEST(AllocatableCUFTest, StreamDeviceAllocatable) {
   cudaDeviceSynchronize();
   EXPECT_EQ(cudaSuccess, cudaGetLastError());
 }
+
+TEST(AllocatableAsyncTest, StreamDeviceAllocatable) {
+  using Fortran::common::TypeCategory;
+  RTNAME(CUFRegisterAllocator)();
+  // REAL(4), DEVICE, ALLOCATABLE :: a(:)
+  auto a{createAllocatable(TypeCategory::Real, 4)};
+  a->SetAllocIdx(kDeviceAllocatorPos);
+  EXPECT_EQ((int)kDeviceAllocatorPos, a->GetAllocIdx());
+  EXPECT_FALSE(a->HasAddendum());
+  RTNAME(AllocatableSetBounds)(*a, 0, 1, 10);
+
+  cudaStream_t stream;
+  cudaStreamCreate(&stream);
+  EXPECT_EQ(cudaSuccess, cudaGetLastError());
+
+  RTNAME(AllocatableAllocate)
+  (*a, /*asyncObject=*/(int64_t *)&stream, /*hasStat=*/false,
+      /*errMsg=*/nullptr, __FILE__, __LINE__);
+  EXPECT_TRUE(a->IsAllocated());
+  cudaDeviceSynchronize();
+  EXPECT_EQ(cudaSuccess, cudaGetLastError());
+  cudaStream_t s = RTDECL(CUFAssociatedGetStream)(a->raw().base_addr);
+  EXPECT_EQ(s, stream);
+  RTNAME(AllocatableDeallocate)
+  (*a, /*hasStat=*/false, /*errMsg=*/nullptr, __FILE__, __LINE__);
+  EXPECT_FALSE(a->IsAllocated());
+  cudaDeviceSynchronize();
+
+  cudaStream_t defaultStream = 0;
+  RTNAME(AllocatableAllocate)
+  (*a, /*asyncObject=*/(int64_t *)&defaultStream, /*hasStat=*/false,
+      /*errMsg=*/nullptr, __FILE__, __LINE__);
+  EXPECT_TRUE(a->IsAllocated());
+  cudaDeviceSynchronize();
+  EXPECT_EQ(cudaSuccess, cudaGetLastError());
+  cudaStream_t d = RTDECL(CUFAssociatedGetStream)(a->raw().base_addr);
+  EXPECT_EQ(d, defaultStream);
+  RTNAME(AllocatableDeallocate)
+  (*a, /*hasStat=*/false, /*errMsg=*/nullptr, __FILE__, __LINE__);
+  EXPECT_FALSE(a->IsAllocated());
+  cudaDeviceSynchronize();
+
+  RTNAME(AllocatableAllocate)
+  (*a, /*asyncObject=*/nullptr, /*hasStat=*/false, /*errMsg=*/nullptr, __FILE__,
+      __LINE__);
+  EXPECT_TRUE(a->IsAllocated());
+  cudaDeviceSynchronize();
+  EXPECT_EQ(cudaSuccess, cudaGetLastError());
+  cudaStream_t empty = RTDECL(CUFAssociatedGetStream)(a->raw().base_addr);
+  EXPECT_EQ(empty, nullptr);
+}

flang/include/flang/Runtime/CUDA/allocator.h

@@ -13,11 +13,14 @@
 #include "flang/Runtime/descriptor-consts.h"
 #include "flang/Runtime/entry-names.h"
 
+#include "cuda_runtime.h"
+
 namespace Fortran::runtime::cuda {
 
 extern "C" {
 
 void RTDECL(CUFRegisterAllocator)();
+cudaStream_t RTDECL(CUFAssociatedGetStream)(void *);
 }
 
 void *CUFAllocPinned(std::size_t, std::int64_t *);

llvm/docs/LangRef.rst

@@ -3234,6 +3234,24 @@ A "convergencectrl" operand bundle is only valid on a ``convergent`` operation.
 When present, the operand bundle must contain exactly one value of token type.
 See the :doc:`ConvergentOperations` document for details.
 
+.. _deactivationsymbol:
+
+Deactivation Symbol Operand Bundles
+^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+
+A ``"deactivation-symbol"`` operand bundle is valid on the following
+instructions (AArch64 only):
+
+- Call to a normal function with ``notail`` attribute and a first argument and
+  return value of type ``ptr``.
+- Call to ``llvm.ptrauth.sign`` or ``llvm.ptrauth.auth`` intrinsics.
+
+This operand bundle specifies that if the deactivation symbol is defined
+to a valid value for the target, the marked instruction will return the
+value of its first argument instead of calling the specified function
+or intrinsic. This is achieved with ``PATCHINST`` relocations on the
+target instructions (see the AArch64 psABI for details).
+
 .. _moduleasm:
 
 Module-Level Inline Assembly
@@ -5284,7 +5302,7 @@ need to refer to the actual function body.
 Pointer Authentication Constants
 --------------------------------
 
-``ptrauth (ptr CST, i32 KEY[, i64 DISC[, ptr ADDRDISC]?]?)``
+``ptrauth (ptr CST, i32 KEY[, i64 DISC[, ptr ADDRDISC[, ptr DS]?]?]?)``
 
 A '``ptrauth``' constant represents a pointer with a cryptographic
 authentication signature embedded into some bits, as described in the
@@ -5313,6 +5331,11 @@ Otherwise, the expression is equivalent to:
     %tmp2 = call i64 @llvm.ptrauth.sign(i64 ptrtoint (ptr CST to i64), i32 KEY, i64 %tmp1)
     %val = inttoptr i64 %tmp2 to ptr
 
+If the deactivation symbol operand ``DS`` has a non-null value,
+the semantics are as if a :ref:`deactivation-symbol operand bundle
+<deactivationsymbol>` were added to the ``llvm.ptrauth.sign`` intrinsic
+calls above, with ``DS`` as the only operand.
+
 .. _constantexprs:
 
 Constant Expressions

llvm/include/llvm/Bitcode/LLVMBitCodes.h

@@ -437,6 +437,8 @@ enum ConstantsCodes {
   CST_CODE_CE_GEP_WITH_INRANGE = 31,  // [opty, flags, range, n x operands]
   CST_CODE_CE_GEP = 32,               // [opty, flags, n x operands]
   CST_CODE_PTRAUTH = 33,              // [ptr, key, disc, addrdisc]
+  CST_CODE_PTRAUTH2 = 34,             // [ptr, key, disc, addrdisc,
+                                      //  deactivation_symbol]
 };
 
 /// CastOpcodes - These are values used in the bitcode files to encode which