This fixes a crash in the attached test case due to missing type
inference for ICmp VPInstructions.

This is similar to c1ad363e, but with
the additional twist that initializing an existing value from a
`MemberExpr` was not working correctly.

If the pattern of a pack indexing type did not contain a pack, we would
still construct a pack indexing type (to improve error messages) but we
would fail to make the type as dependent, leading to infinite recursion
when trying to extract a canonical type.

Fold ((cst << x) & 1) to zext(x == 0) when cst is odd.

Fixes: https://github.com/llvm/llvm-project/issues/73384
Alive2: https://alive2.llvm.org/ce/z/5RbaK6

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>