Summary:
These tests likely always failed but was hidden by the expected return
value. Simply make them require AMDGPU as a registered target so they
don't fail on other machines.

Initialize both elements to 0.

The current interpreter does this, so follow suit to match its
diagnostics.

Summary:
Currently we cannot compile `__builtin_amdgcn_ballot_w64` on non-wave64
targets even though it is valid. This is relevant for making library
code that can handle both without needing to check the wavefront size.
This patch relaxes the semantic check for w64 so it can be used
normally.

Summary:
The PTX language rejects globals with `.` in the name. We need to change
the global name if we are targeting NVPTX to prevent the toolchain from
complaining.

Correct AMDGPU strictfp tests to follow the rules documented in the
LangRef:
https://llvm.org/docs/LangRef.html#constrained-floating-point-intrinsics

These tests needed the strictfp attribute added to function calls and
some declarations.

Some of the tests now pass with D146845, others get farther along and
fail with D146845. The tests revealed that further work is required
in mostly AMDGPU atomics to get the tests passing.

Since I was here anyway I removed the strictfp attribute from some
constrained intrinsic declarations. They have this attribute by default.

Test changes verified with D146845.

[X86] getShuffleComment - use MI description to determine AVX512 masked predicates instead of src index offsets.

This just avoids useless work of adding NoRegister to BaseSet, for
consistency with other places that iterate over all physical registers.

In mixed matmul lowering (e.g., i8 to i32) we're seeing the following
sequence:

  %0 = arith.extsi %src : vector<4x[8]xi8> to vector<4x[8]xi32>
  %1 = vector.extract %0[0] : vector<[8]xi32> from vector<4x[8]xi32>
%lhs = vector.scalable.extract %1[0] : vector<[4]xi32> from
vector<[8]xi32>

  ... (same for rhs)

%2 = vector.outerproduct %lhs, %rhs, %acc vector<[4]xi32>,
vector<[4]xi32>

  // x4 chained by accumulator

This chain of 4 outer products can be fused into a single 4-way widening
variant but the pass doesn't match on the IR, as it expects the source
of the inputs to be an extend and it can't look through the extracts.

This patch fixes this with two rewrites that swaps extract(extend) into
extend(extract).

Related to #78975, #79288.

The "Refer to" and table shouldn't be in the example code sequence.

Fixes #77659
Fixes #46357

Picked up from https://reviews.llvm.org/D114119

This adds the support to add the target-feature to outline atomic operations (calling the
runtime library instead).

[X86] addConstantComments - split VPERMILPS/VPERMILPD handling to reduce repeated switch cases etc. NFC.

This is to add test coverage for crash report in #80340

This enables IR expansion for i128 divisions. The vector case is still
broken because ExpandLargeDivRem doesn't try to handle them.

Fixes: SWDEV-426193

Fixes #80366

Implement PhiLoweringHelper for GlobalISel in DivergenceLoweringHelper.
Use machine uniformity analysis to find divergent i1 phis and select
them as lane mask phis in same way SILowerI1Copies select VReg_1 phis.
Note that divergent i1 phis include phis created by LCSSA and all cases
of uses outside of cycle are actually covered by "lowering LCSSA phis".
GlobalISel lane masks are registers with sgpr register class and S1 LLT.

TODO: General goal is that instructions created in this pass are fully
instruction-selected so that selection of lane mask phis is not split
across multiple passes.

patch 3 from: https://github.com/llvm/llvm-project/pull/73337

The SGPR registers used for preserving EXEC mask while lowering the
whole-wave register spills and copies should be preserved at the prolog
and epilog if they are in the CSR range. It isn't happening when there
is only wwm-copy lowered and there are no wwm-spills. This patch
addresses that problem.

This patch adds support for forwarding the target-cpu and
target-features attributes to functions outlined in the OpenMPIRBuilder.
This, in turn, results in the addition of these attributes for functions
created during the translation of the `omp.parallel`, `omp.task` and
`omp.teams` operations, and for the `omp.wsloop` operation when doing
codegen for an OpenMP target device.

[X86] X86FixupVectorConstants - load+zero vector constants that can be stored in a truncated form (#80428)

Further develops the vsextload support added in #79815 / b5d35fea - reduces the size of the vector constant by storing it in the constant pool in a truncated form, and zero-extend it as part of the load.

Fix: https://github.com/llvm/llvm-project/issues/78568



---------

Co-authored-by: jeanPerier <jean.perier.polytechnique@gmail.com>

This is a separate, but related issue to #69152 that was attempting to improve
AA with scalable dependency distances. This patch attempts to improve when
there are scalable accesses with a constant offset between them. We happen to
get a report of such a thing recently, where so long as the vscale_range is
known, the maximum size of the access can be assessed and better aliasing
results can be returned.

The Upper range of the vscale_range, along with known part of the typesize are
used to prove that Off >= CR.upper * LSize. It does not try to produce
PartialAlias results at the moment from the lower vscale_range. It also enables
the added benefit of allowing better alias analysis when the RHS of the two
values is scalable, but the LHS is normal and can be treated like any other
aliasing query.

Even when `private-function-dynamic-ownership` is set, ownership should
never be passed to the callee. This can lead to double deallocs (#77096)
or use-after-free in the caller because ownership is currently passed
regardless of whether there are any further uses of the buffer in the
caller or not.

Note: This is consistent with the fact that ownership is never passed to
nested regions.

This commit fixes #77096.