Consistently add `branch_weights` metadata in any condition branch
created by `LoopVectorize.cpp`:
- Will only add metadata if the original loop-latch branch had metadata
assigned.
- Most checks should rarely trigger so I am using a 127:1 ratio.
- For the middle block we assume an equal distribution of modulo
results.

Add initial support for DIExpression in LLVM dialect.

Similar to LLVM IR, DI Expression is encoded as a list of uint64. The
difference is that LLVM IR has helpers for understanding the expression
(e.g. for verification and pretty printing), whereas the current support
added by this PR treats the expression elements as opaque.

This is an alternative to #71420 that handles i32 on RV64 safely
by pre-promoting the pattern in DAG combine.

We can use (ADDI (BSET X0, X), -1).

This patch fixes:

  clang/lib/StaticAnalyzer/Checkers/StdLibraryFunctionsChecker.cpp:609:23:
  error: 'describe' overrides a member function but is not marked
  'override' [-Werror,-Winconsistent-missing-override]

  clang/lib/StaticAnalyzer/Checkers/StdLibraryFunctionsChecker.cpp:627:23:
  error: 'describe' overrides a member function but is not marked
  'override' [-Werror,-Winconsistent-missing-override]

Enable -rpath, -shared, and -static for the flang frontend. This brings
it in line with clang. Fixes issue #65546.

Add `setBranchWeights` convenience function to ProfDataUtils.h and use
it where appropriate.

This reverts commit 0718c1a8.

The test doesn't have the correct requires string apparently because it
attempts to run on all the buildbots. Reverting until I have time to fix
the test and reland it.

This patch makes it so that specifying an invalid value for
CMAKE_BUILD_TYPE is a fatal error. Having this simply as a warning has
caused me (and probably others) a decent amount of headache. The check
was present before, but was proposed to be modified to a warning in
https://github.com/llvm/llvm-project/issues/60975 and changed to a
warning in c75dbeda. This patch
reenables that behavior to hopefully reduce frustration for people
building LLVM in the common case while still allowing for alternative
build types to be setup without needing to perform source modification
through the addition of a CMake flag.

I made a typo at some point while doing the subprocess implementation,
trying to pop RIP from the stack. For whatever reason, this ends up as
popq %rax after being JITed, which means we aren't properly preserving
the value of %rdi.

Regression test added.

This fixes #72188.

This change adds the image used by the self-hosted Github Actions
builders.

In an attempt to make the transition simple, all of the different images
share as much of the same state as possible, including packages, users,
etc... This results in bigger images, but that shouldn't be a problem.
That said, the refactorings caused the buildkite image to shrink by 100
MB.

This change also renames all of the packages for consistency. Bots will
have to be changed to use the new package names eventually.

Again, docker-compose was used as the source of truth for defining
argument.

I have already pushed example images to ghcr.io/libcxx/<name>:testing

Fixes #72433

Rollforward of #69032, which was reverted in
[63d19cfd](https://github.com/llvm/llvm-project/commit/63d19cfd853b0b4e374f40842ee388b0da5de36f).

New: implemented changes in
https://github.com/llvm/llvm-project/pull/69032#issuecomment-1809250162.
Given the PassBuilder is how we expect users to register passes, the
tutorial should reflect that.

This reverts commit 13fe0386.

May have broken an LLDB test https://lab.llvm.org/buildbot/#/builders/96/builds/48609

The checker now displays one combined note tag for errno-related and
"case"-related notes. Previous functions in the errno-modeling part that
were used for construction of note tags are removed. The note tag added
by StdLibraryFunctionsChecker contains the code to display the note tag
for 'errno' (this was done previously by these removed functions).

Let `llvm.launder.invariant.group` intrinsic as well as instructions
operating on memory addresses, whose invariance may be broken by the
intrinsic, to be rewritten.

Fixes: https://github.com/llvm/llvm-project/issues/72035.

Both options do not affect the AST content that is serialized into the PCM. This
commit includes the following changes:
    
1.) Mark `-fvisibility={}` and `-ftype-visibility={}` as benign options.That
     means they are no longer considered part of the module hash, which can
     reduce the number of module variants.
    
2.) Add a test to verify the generated LLVM IR is not affected by the default
     visibiliy mode in the module.

3.) Add a test to clang-scan-deps to ensure only one module is build, even if
      the above mentioned options are used.
    
This fixes rdar://118246054.

This patch adds initial support to the NVPTX target, enabling `flang` to
produce OpenMP offload code for NVPTX targets.

models, NFC.

For the following testcase:

undef %1.sub1:sgpr_96 = COPY undef %0:sgpr_32
%3:vgpr_32 = V_LSHL_ADD_U32_e64 %1.sub1:sgpr_96, ...

GCNRewritePartialRegUses produced:

%4:vgpr_32 = COPY undef %1:sgpr_32
dead %2:vgpr_32 = V_LSHL_ADD_U32_e64 %4, ...

Register class for %4 is incorrect: there should be sgpr_32 instead of
vgpr_32 because the original %1 had scalar regclass. This patch fixes
that.

Note that GCNRewritePartialRegUses pass isn't enabled by default yet.

Need to check if we ran into subvector extract pattern before checking
for identity vector to avoid compiler crash.

This introduces the Armv9.5-A architecture version, including the
relevant command-line option for -march.

Mode details about the Armv9.5-A architecture version can be found at:
*
https://community.arm.com/arm-community-blogs/b/architectures-and-processors-blog/posts/arm-a-profile-architecture-developments-2023
* https://developer.arm.com/documentation/ddi0602/2023-09/

Patch by Oliver Stannard.

Opaque ptr cleanup effort (NFC).

operand.

Need to copy the submask not to the very first part of the common
extractelements vector mask, but to the proper one to avoid wrong code
emission.

None of the in-tree implementations have different behaviour for
physical vs virtual registers, and it seems would work equally well if
used with virtual registers. As such, perhaps it's simplest to just drop
that part of the doc comment.

This hook is called by the target-independent implementation of
TargetInstrInfo::describeLoadedValue. I've opted to test it via a C++
unit test, which although fiddly to set up seems the right way to test a
function with such clear intended semantics (rather than testing the
impact indirectly).

isAddImmediate will never recognise ADDIW as an add immediate which I
_think_ is conservatively correct, as the caller may not understand its
semantics vs ADDI.

Note that although the doc comment for isAddImmediate specifies its
behaviour solely in terms of physical registers, none of the current
in-tree implementations (including this one) bail out on virtual
registers (see #72357).

vectorized extractelements operands.

Enable this flow by -riscv-split-regalloc=1 (default disable), and could
designate specific allocator to RVV by
-riscv-rvv-regalloc=<fast|basic|greedy>

It uses the RegClass filter function to decide which regclass need to be
processed.

This patch is pre-requirement for supporting PostRA vsetvl insertion
pass.

I missed this test in PR https://github.com/llvm/llvm-project/pull/72304

Adds the following SME2 builtins:

svminnm_single_f(16|32|64)_x(2|4)
svminnm_f(16|32|64)_x(2|4)
svmaxnm_single_f(16|32|64)_x(2|4)
svmaxnm_f(16|32|64)_x(2|4)

See
[ARM-software/acle#217](https://github.com/ARM-software/acle/pull/217)

Patch by: Kerry McLaughlin <kerry.mclaughlin@arm.com>

Make trivial ctor/dtor implicitly host device functions so that they can
be used to initialize file-scope
device variables to match nvcc behavior.

Fixes: https://github.com/llvm/llvm-project/issues/72261

Fixes: SWDEV-432412

While the `gpu.launch` Op allows setting the size via the
`dynamic_shared_memory_size` argument, accessing the dynamic shared
memory is very convoluted. This PR implements the proposed Op,
`gpu.dynamic_shared_memory` that aims to simplify the utilization of
dynamic shared memory.

RFC:
https://discourse.llvm.org/t/rfc-simplifying-dynamic-shared-memory-access-in-gpu/

**Proposal from RFC**
This PR `gpu.dynamic.shared.memory` Op to use dynamic shared memory
feature efficiently. It is is a powerful feature that enables the
allocation of shared memory at runtime with the kernel launch on the
host. Afterwards, the memory can be accessed directly from the device. I
believe similar story exists for AMDGPU.

**Current way Using Dynamic Shared Memory with MLIR**

Let me illustrate the challenges of using dynamic shared memory in MLIR
with an example below. The process involves several steps:
- memref.global 0-sized array LLVM's NVPTX backend expects
- dynamic_shared_memory_size Set the size of dynamic shared memory
- memref.get_global Access the global symbol
- reinterpret_cast and subview Many OPs for pointer arithmetic

```
// Step 1. Create 0-sized global symbol. Manually set the alignment
memref.global "private" @dynamicShmem  : memref<0xf16, 3> { alignment = 16 }
func.func @main() {
  // Step 2. Allocate shared memory
  gpu.launch blocks(...) threads(...)
    dynamic_shared_memory_size %c10000 {
    // Step 3. Access the global object
    %shmem = memref.get_global @dynamicShmem : memref<0xf16, 3>
    // Step 4. A sequence of `memref.reinterpret_cast` and `memref.subview` operations.
    %4 = memref.reinterpret_cast %shmem to offset: [0], sizes: [14, 64, 128],  strides: [8192,128,1] : memref<0xf16, 3> to memref<14x64x128xf16,3>
    %5 = memref.subview %4[7, 0, 0][7, 64, 128][1,1,1] : memref<14x64x128xf16,3> to memref<7x64x128xf16, strided<[8192, 128, 1], offset: 57344>, 3>
    %6 = memref.subview %5[2, 0, 0][1, 64, 128][1,1,1] : memref<7x64x128xf16, strided<[8192, 128, 1], offset: 57344>, 3> to memref<64x128xf16, strided<[128, 1], offset: 73728>, 3>
    %7 = memref.subview %6[0, 0][64, 64][1,1]  : memref<64x128xf16, strided<[128, 1], offset: 73728>, 3> to memref<64x64xf16, strided<[128, 1], offset: 73728>, 3>
    %8 = memref.subview %6[32, 0][64, 64][1,1] : memref<64x128xf16, strided<[128, 1], offset: 73728>, 3> to memref<64x64xf16, strided<[128, 1], offset: 77824>, 3>
    // Step.5 Use
    "test.use.shared.memory"(%7) : (memref<64x64xf16, strided<[128, 1], offset: 73728>, 3>) -> (index)
    "test.use.shared.memory"(%8) : (memref<64x64xf16, strided<[128, 1], offset: 77824>, 3>) -> (index)
    gpu.terminator
  }
```

Let’s write the program above with that:

```
func.func @main() {
    gpu.launch blocks(...) threads(...) dynamic_shared_memory_size %c10000 {
    	%i = arith.constant 18 : index
        // Step 1: Obtain shared memory directly
        %shmem = gpu.dynamic_shared_memory : memref<?xi8, 3>
        %c147456 = arith.constant 147456 : index
        %c155648 = arith.constant 155648 : index
        %7 = memref.view %shmem[%c147456][] : memref<?xi8, 3> to memref<64x64xf16, 3>
        %8 = memref.view %shmem[%c155648][] : memref<?xi8, 3> to memref<64x64xf16, 3>

        // Step 2: Utilize the shared memory
        "test.use.shared.memory"(%7) : (memref<64x64xf16, 3>) -> (index)
        "test.use.shared.memory"(%8) : (memref<64x64xf16, 3>) -> (index)
    }
}
```

This PR resolves #72513

This PR introduce `cp.async.bulk.tensor.shared.cluster.global.multicast`
Op in NVVM dialect. It loads data using TMA data from global memory to
shared memory of multiple CTAs in the cluster.

It resolves #72368

Hoist non-atomic expressions in atomic intrinsics. Use a list to collect
the non-atomic expressions since the max and min intrinsics can have
more than two operands.

Hoisting makes the lowering to LLVMIR of iand,ior,ieor intrinsics
trivial. For max and min this still results in multiple instructions in
the atomic region but the loads are removed, which should help improve
performance.