Merge branch 'llvm:main' into fix-bindc-name

Author: eugeneepshteyn

bolt/test/runtime/copy_file.py

@@ -0,0 +1,15 @@
+import sys
+import shutil
+
+with open(sys.argv[1] + ".output") as log_file:
+    lines = log_file.readlines()
+    for line in lines:
+        if line.startswith(sys.argv[2]):
+            pid = line.split(" ")[1].strip()
+            shutil.copy(
+                sys.argv[1] + "." + pid + ".fdata",
+                sys.argv[1] + "." + sys.argv[3] + ".fdata",
+            )
+            sys.exit(0)
+
+sys.exit(1)

bolt/test/runtime/instrumentation-indirect-2.c

@@ -50,7 +50,7 @@ int main() {
   return 0;
 }
 /*
-REQUIRES: system-linux,shell,fuser
+REQUIRES: system-linux,fuser
 
 RUN: %clang %cflags %s -o %t.exe -Wl,-q -pie -fpie
 
@@ -61,10 +61,14 @@ RUN: --instrumentation-wait-forks
 
 # Instrumented program needs to finish returning zero
 # Both output and profile must contain all 16 functions
-RUN: %t.instrumented_conservative > %t.output
-# Wait for profile and output to be fully written
-RUN: bash %S/wait_file.sh %t.output
-RUN: bash %S/wait_file.sh %t.fdata
+# We need to use bash to invoke this as otherwise we hang inside a
+# popen.communicate call in lit's internal shell. Eventually we should not
+# need this.
+# TODO(boomanaiden154): Remove once
+# https://github.com/llvm/llvm-project/issues/156484 is fixed.
+RUN: bash -c "%t.instrumented_conservative; wait" > %t.output
+# We can just read because we ensure the profile will be fully written by
+# calling wait inside the bash invocation.
 RUN: cat %t.output | FileCheck %s --check-prefix=CHECK-OUTPUT
 RUN: cat %t.fdata | FileCheck %s --check-prefix=CHECK-COMMON-PROF
 
@@ -112,14 +116,8 @@ RUN: bash %S/wait_file.sh %t.output
 # Make sure all functions were called
 RUN: cat %t.output | FileCheck %s --check-prefix=CHECK-OUTPUT
 
-RUN: child_pid=$(cat %t.output | grep funcA | awk '{print $2;}')
-RUN: par_pid=$(cat %t.output | grep funcB | awk '{print $2;}')
-
-RUN: bash %S/wait_file.sh %t.$child_pid.fdata
-RUN: bash %S/wait_file.sh %t.$par_pid.fdata
-
-RUN: mv %t.$child_pid.fdata %t.child.fdata
-RUN: mv %t.$par_pid.fdata %t.parent.fdata
+RUN: %python %S/copy_file.py %t funcA child
+RUN: %python %S/copy_file.py %t funcB parent
 
 # Instrumented binary must produce two profiles with only local calls
 # recorded. Functions called only in child should not appear in parent's

clang/include/clang/Basic/BuiltinsAMDGPU.def

@@ -667,6 +667,9 @@ TARGET_BUILTIN(__builtin_amdgcn_global_load_monitor_b128, "V4iV4i*1Ii", "nc", "g
 TARGET_BUILTIN(__builtin_amdgcn_flat_load_monitor_b32, "ii*0Ii", "nc", "gfx1250-insts")
 TARGET_BUILTIN(__builtin_amdgcn_flat_load_monitor_b64, "V2iV2i*0Ii", "nc", "gfx1250-insts")
 TARGET_BUILTIN(__builtin_amdgcn_flat_load_monitor_b128, "V4iV4i*0Ii", "nc", "gfx1250-insts")
+TARGET_BUILTIN(__builtin_amdgcn_cluster_load_b32, "ii*1Iii", "nc", "gfx1250-insts,wavefrontsize32")
+TARGET_BUILTIN(__builtin_amdgcn_cluster_load_b64, "V2iV2i*1Iii", "nc", "gfx1250-insts,wavefrontsize32")
+TARGET_BUILTIN(__builtin_amdgcn_cluster_load_b128, "V4iV4i*1Iii", "nc", "gfx1250-insts,wavefrontsize32")
 TARGET_BUILTIN(__builtin_amdgcn_global_load_async_to_lds_b8, "vc*1c*3IiIi", "nc", "gfx1250-insts")
 TARGET_BUILTIN(__builtin_amdgcn_global_load_async_to_lds_b32, "vi*1i*3IiIi", "nc", "gfx1250-insts")
 TARGET_BUILTIN(__builtin_amdgcn_global_load_async_to_lds_b64, "vV2i*1V2i*3IiIi", "nc", "gfx1250-insts")

clang/include/clang/Basic/DiagnosticDriverKinds.td

@@ -361,9 +361,6 @@ def err_drv_optimization_remark_format : Error<
 def err_drv_no_neon_modifier : Error<"[no]neon is not accepted as modifier, please use [no]simd instead">;
 def err_drv_invalid_omp_target : Error<"OpenMP target is invalid: '%0'">;
 def err_drv_incompatible_omp_arch : Error<"OpenMP target architecture '%0' pointer size is incompatible with host '%1'">;
-def err_drv_omp_host_ir_file_not_found : Error<
-  "provided host compiler IR file '%0' is required to generate code for OpenMP "
-  "target regions but cannot be found">;
 def err_drv_omp_host_target_not_supported : Error<
   "target '%0' is not a supported OpenMP host target">;
 def err_drv_expecting_fopenmp_with_fopenmp_targets : Error<

clang/include/clang/Basic/DiagnosticFrontendKinds.td

@@ -312,6 +312,9 @@ def err_target_unsupported_type_for_abi
     : Error<"%0 requires %1 type support, but ABI '%2' does not support it">;
 }
 
+def err_omp_host_ir_file_not_found : Error<
+  "provided host compiler IR file '%0' is required to generate code for OpenMP "
+  "target regions but cannot be found">;
 def err_alias_to_undefined : Error<
   "%select{alias|ifunc}0 must point to a defined "
   "%select{variable or |}1function">;

clang/include/clang/CIR/Dialect/Builder/CIRBaseBuilder.h

@@ -160,16 +160,15 @@ class CIRBaseBuilderTy : public mlir::OpBuilder {
   }
 
   cir::LoadOp createLoad(mlir::Location loc, mlir::Value ptr,
-                         uint64_t alignment = 0) {
+                         bool isVolatile = false, uint64_t alignment = 0) {
     mlir::IntegerAttr alignmentAttr = getAlignmentAttr(alignment);
-    assert(!cir::MissingFeatures::opLoadStoreVolatile());
-    return cir::LoadOp::create(*this, loc, ptr, /*isDeref=*/false,
+    return cir::LoadOp::create(*this, loc, ptr, /*isDeref=*/false, isVolatile,
                                alignmentAttr, cir::MemOrderAttr{});
   }
 
   mlir::Value createAlignedLoad(mlir::Location loc, mlir::Value ptr,
                                 uint64_t alignment) {
-    return createLoad(loc, ptr, alignment);
+    return createLoad(loc, ptr, /*isVolatile=*/false, alignment);
   }
 
   mlir::Value createNot(mlir::Value value) {
@@ -250,7 +249,7 @@ class CIRBaseBuilderTy : public mlir::OpBuilder {
                            bool isVolatile = false,
                            mlir::IntegerAttr align = {},
                            cir::MemOrderAttr order = {}) {
-    return cir::StoreOp::create(*this, loc, val, dst, align, order);
+    return cir::StoreOp::create(*this, loc, val, dst, isVolatile, align, order);
   }
 
   [[nodiscard]] cir::GlobalOp createGlobal(mlir::ModuleOp mlirModule,
@@ -274,7 +273,8 @@ class CIRBaseBuilderTy : public mlir::OpBuilder {
     mlir::IntegerAttr alignmentAttr = getAlignmentAttr(alignment);
     auto addr = createAlloca(loc, getPointerTo(type), type, {}, alignmentAttr);
     return cir::LoadOp::create(*this, loc, addr, /*isDeref=*/false,
-                               alignmentAttr, /*mem_order=*/{});
+                               /*isVolatile=*/false, alignmentAttr,
+                               /*mem_order=*/{});
   }
 
   cir::PtrStrideOp createPtrStride(mlir::Location loc, mlir::Value base,

clang/include/clang/CIR/Dialect/IR/CIROps.td

@@ -404,7 +404,11 @@ def CIR_LoadOp : CIR_Op<"load", [
     `cir.load` reads a value (lvalue to rvalue conversion) given an address
     backed up by a `cir.ptr` type. A unit attribute `deref` can be used to
     mark the resulting value as used by another operation to dereference
-    a pointer.
+    a pointer. A unit attribute `volatile` can be used to indicate a volatile
+    loading. Load can be marked atomic by using `atomic(<mem_order>)`.
+
+    `alignment` can be used to specify an alignment that's different from the
+    default, which is computed from `result`'s type ABI data layout.
 
     Example:
 
@@ -416,18 +420,26 @@ def CIR_LoadOp : CIR_Op<"load", [
     // Load address from memory at address %0. %3 is used by at least one
     // operation that dereferences a pointer.
     %3 = cir.load deref %0 : !cir.ptr<!cir.ptr<i32>>
+
+    // Perform a volatile load from address in %0.
+    %4 = cir.load volatile %0 : !cir.ptr<i32>, i32
+
+    // Others
+    %x = cir.load align(16) atomic(seq_cst) %0 : !cir.ptr<i32>, i32
     ```
   }];
 
   let arguments = (ins Arg<CIR_PointerType, "the address to load from",
                            [MemRead]>:$addr,
                        UnitAttr:$isDeref,
+                       UnitAttr:$is_volatile,
                        OptionalAttr<I64Attr>:$alignment,
                        OptionalAttr<CIR_MemOrder>:$mem_order);
   let results = (outs CIR_AnyType:$result);
 
   let assemblyFormat = [{
     (`deref` $isDeref^)?
+    (`volatile` $is_volatile^)?
     (`align` `(` $alignment^ `)`)?
     (`atomic` `(` $mem_order^ `)`)?
     $addr `:` qualified(type($addr)) `,` type($result) attr-dict
@@ -452,24 +464,32 @@ def CIR_StoreOp : CIR_Op<"store", [
     a volatile store. Store's can be marked atomic by using
     `atomic(<mem_order>)`.
 
-    `align` can be used to specify an alignment that's different from the
+    `alignment` can be used to specify an alignment that's different from the
     default, which is computed from `result`'s type ABI data layout.
 
     Example:
 
     ```mlir
     // Store a function argument to local storage, address in %0.
     cir.store %arg0, %0 : i32, !cir.ptr<i32>
+
+    // Perform a volatile store into memory location at the address in %0.
+    cir.store volatile %arg0, %0 : i32, !cir.ptr<i32>
+
+    // Others
+    cir.store align(16) atomic(seq_cst) %x, %addr : i32, !cir.ptr<i32>
     ```
   }];
 
   let arguments = (ins CIR_AnyType:$value,
                        Arg<CIR_PointerType, "the address to store the value",
                            [MemWrite]>:$addr,
+                       UnitAttr:$is_volatile,
                        OptionalAttr<I64Attr>:$alignment,
                        OptionalAttr<CIR_MemOrder>:$mem_order);
 
   let assemblyFormat = [{
+    (`volatile` $is_volatile^)?
     (`align` `(` $alignment^ `)`)?
     (`atomic` `(` $mem_order^ `)`)?
     $value `,` $addr attr-dict `:` type($value) `,` qualified(type($addr))

clang/include/clang/CIR/MissingFeatures.h

@@ -47,9 +47,8 @@ struct MissingFeatures {
   // Load/store attributes
   static bool opLoadStoreThreadLocal() { return false; }
   static bool opLoadEmitScalarRangeCheck() { return false; }
-  static bool opLoadBooleanRepresentation() { return false; }
+  static bool opLoadStoreNontemporal() { return false; }
   static bool opLoadStoreTbaa() { return false; }
-  static bool opLoadStoreVolatile() { return false; }
   static bool opLoadStoreAtomic() { return false; }
   static bool opLoadStoreObjC() { return false; }
 

clang/include/clang/Driver/Options.td

@@ -8719,7 +8719,8 @@ def fopenmp_is_target_device : Flag<["-"], "fopenmp-is-target-device">,
   HelpText<"Generate code only for an OpenMP target device.">;
 def : Flag<["-"], "fopenmp-is-device">, Alias<fopenmp_is_target_device>;
 def fopenmp_host_ir_file_path : Separate<["-"], "fopenmp-host-ir-file-path">,
-  HelpText<"Path to the IR file produced by the frontend for the host.">;
+  HelpText<"Path to the IR file produced by the frontend for the host.">,
+  MarshallingInfoString<LangOpts<"OMPHostIRFile">>;
 
 } // let Visibility = [CC1Option, FC1Option]
 

clang/lib/AST/ByteCode/Pointer.cpp

@@ -627,9 +627,6 @@ bool Pointer::hasSameBase(const Pointer &A, const Pointer &B) {
   if (A.isTypeidPointer() && B.isTypeidPointer())
     return true;
 
-  if (A.isIntegralPointer() || B.isIntegralPointer())
-    return A.getSource() == B.getSource();
-
   if (A.StorageKind != B.StorageKind)
     return false;