BSD Commit Log Search

[AMDGPU] Allow hazard checks for WMMA co-exec

Now we are just inserting V_NOP instrtuctions, try to schedule
something into the shadow.

It is still somewhat imprecise, for example AdvanceCycle() will
use TII.getNumWaitStates() anyway, but in a scheduling mode
we are not required to be precise. We must be finally precise
in the hazard recognizer mode. Then EmittedInstrs buffer is also
limited to MaxLookAhead even though VALU only hazards may actually
never expire and require an endless buffer. But that's OK, we can
at least mitigate what the buffer can hold. The buffer is also
currently much bigger than any of VALU hazards may need.

That said the rest of the 'fix*' functions here can be changed
the same way, these which are using V_NOPs. This one is just the
worst because it may require up to 9 nops.

[InstSimplify] Extend icmp-of-add simplification to sle/sgt/sge (#168900)

When comparing additions with the same base where one has `nsw`, the
following simplification can be performed:

```llvm
icmp slt/sgt/sle/sge (x + C1), (x +nsw C2)
=>
icmp slt/sgt/sle/sge C1, C2
```

Previously this was only done for `slt`. This patch extends it to the
`sgt`, `sle`, and `sge` predicates when either of the conditions hold:
- `C1 <= C2 && C1 >= 0`, or
- `C2 <= C1 && C1 <= 0`

This patch also handles the `C1 == C2` case, which was previously
excluded.

Proof: https://alive2.llvm.org/ce/z/LtmY4f

[compiler-rt] [test] Generalize an UNSUPPORTED marking (#168858)

Don't specifically target windows-msvc - the same goes for any windows
target; mingw doesn't have dlfcn.h either.

[AMDGPU] Precommit test for issue in amdgpu-rewrite-agpr-copy-mfma, (#168609)

which reassigns scale operand in vgpr_32 register to agpr_32, not
permitted by instruction format. Reduced from ck.

---------

Co-authored-by: Matt Arsenault <arsenm2 at gmail.com>
Co-authored-by: theRonShark <ron.lieberman at amd.com>

[LV] Add test a low-trip count test without folding the tail.

Add a low trip count test that is currently vectorized but unprofitable,
for https://github.com/llvm/llvm-project/issues/167858.

[msan] Fix handling of 256-bit hadd/hsub instructions (#168121)

These horizontal add/sub instructions are currently handled by
adding/subtracting tuples of the first operand, followed by tuples of
the second operand. This is not the correct semantics for the 256-bit
insructions: they process the first half of the first operand, then the
first half of the second operand, then the second half of the first
operand, and finally the second half of the second operand (trust me bro
[*]).

This patch fixes the issue by applying the "shards" functionality that
was added in https://github.com/llvm/llvm-project/pull/167954, to handle
the top and bottom 128-bit "shards" in turn.

[*] clang/test/CodeGen/X86/avx2-builtins.c:
```
TEST_CONSTEXPR(match_v8si(_mm256_hadd_epi32(
    (__m256i)(__v8si){10, 20, 30, 40, 50, 60, 70, 80},
    (__m256i)(__v8si){5, 15, 25, 35, 45, 55, 65, 75}),
    30,70,20,60,110,150,100,140));
```

[flang][cuda] Extract element count computation into helper function (#168937)

This patch extracts the common logic for computing array element counts
from shape operands into a reusable helper function in CUFCommon.

[JITLINK] Disable ELF_ppc64_relocations.s on SystemZ host (#168939)

Mark ELF_ppc64_relocations.s as unsupported on SystemZ because of cross
build issue related to using dlsym for host symbols.
Test fails to resolve __tls_get_aadr on SystemZ host.

Co-authored-by: anoopkg6 <anoopkg6 at github.com>

[libc] Removed unused flags from baremetal cache files (#168942)

These flags are not needed for building libc.

[Utils][update_llc_test_checks] Support generating asm tests from templates.

Reduces the pain of manual editing tests applying the same
changes over multiple instructions and keeping them consistent.

[LangRef] Docs: more detailed categories for Vector intrinsics (#168924)

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

[bazel] Bump to 8.4.2 (#168933)

Just staying up to date

Add 64-bit target check for signed AMO builtins

[mlir][llvm] Handle debug record import edge cases

This commit enables the direct import of debug records by default and
fixes issues with two edge cases:
- Detect early on if the address operand is an argument list
  (calling getAddress() for argument lists asserts)
- Use getAddress() to check if the address operand is null, which
  means the address operand is an empty metadata node, which currently
  is not supported.
- Add support for debug label records.

This is a follow-up to:
https://github.com/llvm/llvm-project/pull/167812

DAG: Handle poison in m_Undef

[bazel] Add explicit dep on protobuf (#168928)

This is required for correctly loading the protobuf rules. It's
possible we could drop the version here to a lower version, as long as
that version supports the versions of bazel we support. I picked this
because it is the current version being used by bazel 8.0.0 (which is
defined in the .bazelversion). Users can override this in their project
anyways if they need an older one

Skeleton

[OpenMP] Introduce "loop sequence" as directive association

OpenMP 6.0 introduced a `fuse` directive, and with it a "loop sequence"
as the associated code. What used to be "loop association" has become
"loop-nest association".

Rename Association::Loop to LoopNest, add Association::LoopSeq to
represent the "loop sequence" association.

Change the association of fuse from "block" to "loop sequence".

[clang-doc][NFC] Remove unused headers (#168806)

Removes unused headers or replaces them with headers that directly
provide the symbol instead. For example, `Serialize.h` included `AST.h`,
but it was actually `Serialize.cpp` that needed concept expressions, so
now it includes just `ExprConcepts.h`.

AMDGPU: Use ConstantPool as source value for DAG lowered kernarg loads

This isn't quite a constant pool, but probably close enough for this
purpose. We just need some known invariant value address. The aliasing
queries against the real kernarg base pointer will falsely report
no aliasing, but for invariant memory it probably doesn't matter.

AMDGPU: Handle invariant when lowering global loads

Global with invariant should be treated identically to
constant.

[VPlan] Remove PtrIV::IsScalarAfterVectorization, use VPlan analysis. (#168289)

Remove `VPWidenPointerInductionRecipe::IsScalarAfterVectorization` and
replace it with `onlyScalarValuesUsed`. This removes the need to carry
state from the legacy cost model through VPlan, and the VPlan-based
analysis gives more accurate results, avoiding a number of extracts.

PR: https://github.com/llvm/llvm-project/pull/168289

[Support] Add vector::erase to JSON::Array (#168835)

AMDGPU: Add baseline test for split/widen invariant loads

[ASAN] Disable broken __SANITIZER_DISABLE_CONTAINER_OVERFLOW__ tests on iOS/Android (#168821)

The tests added by #163468 appear to be broken due to lack of libcxx support (?).

Marking unsupported everywhere for now since it passes on some platforms and fails on others, and
I don't know the full list.

Android fail: https://lab.llvm.org/buildbot/#/builders/186/builds/14106

correct logic

Created using spr 1.3.7

[SLP]Check if the non-schedulable phi parent node has unique operands

Need to check if the non-schedulable phi parent node has unique
operands, if the incoming node has copyables, and the node is
commutative. Otherwise, there might be issues with the correct
calculation of the dependencies.

Fixes #168589

[bazel] Delete WORKSPACE file (#168926)

This has been replaced by the MODULE.bazel file. Users can still use
their own WORKSPACE files, but they didn't inherit this file anyways.
Users should migrate to bzlmod as with bazel 9.x that is required.

undo stray change

Created using spr 1.3.7

[LV] Add tests for loops with low trip counts requiring tail-folding.

Add extra tests for over-eager tail-folding for tiny trip-count loops.

Reduced from https://github.com/llvm/llvm-project/issues/167858.