llvm-project.git/llvm/test/Transforms/FunctionAttrs, branch users/nico/python-2

[FunctionAttrs] Handle ConstantRange overflow in memset initializes inference (#145739)

2025-07-01T13:04:52+00:00

Avoid constructing invalid ConstantRange when Offset + Length in memset
overflows signed 64-bit integer space. This prevents assertion failures
when inferring the initializes attribute.

Fixes #140345

[FunctionAttrs][IR] Fix memory attr inference for volatile mem intrinsics (#122926)

2025-06-25T07:29:37+00:00

Per LangRef volatile operations can read and write inaccessible memory:

> any volatile operation can read and/or modify state which is not
> accessible via a regular load or store in this module

Model this by adding inaccessible memory effects in getMemoryEffects()
if the operation is volatile.

In the future, we should model volatile using operand bundles instead.

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

Reland "[NVPTX] Unify and extend barrier{.cta} intrinsic support" (#141143)

2025-05-23T02:38:10+00:00

Note: This relands #140615 adding a ".count" suffix to the non-".all"
variants.

Our current intrinsic support for barrier intrinsics is confusing and
incomplete, with multiple intrinsics mapping to the same instruction and
intrinsic names not clearly conveying intrinsic semantics. Further, we
lack support for some variants. This change unifies the IR
representation to a single consistently named set of intrinsics.

- llvm.nvvm.barrier.cta.sync.aligned.all(i32)
- llvm.nvvm.barrier.cta.sync.aligned.count(i32, i32)
- llvm.nvvm.barrier.cta.arrive.aligned.count(i32, i32)
- llvm.nvvm.barrier.cta.sync.all(i32)
- llvm.nvvm.barrier.cta.sync.count(i32, i32)
- llvm.nvvm.barrier.cta.arrive.count(i32, i32)

The following Auto-Upgrade rules are used to maintain compatibility with
IR using the legacy intrinsics:

* llvm.nvvm.barrier0 --> llvm.nvvm.barrier.cta.sync.aligned.all(0)
* llvm.nvvm.barrier.n --> llvm.nvvm.barrier.cta.sync.aligned.all(x)
* llvm.nvvm.bar.sync --> llvm.nvvm.barrier.cta.sync.aligned.all(x)
* llvm.nvvm.barrier --> llvm.nvvm.barrier.cta.sync.aligned.count(x, y)
* llvm.nvvm.barrier.sync --> llvm.nvvm.barrier.cta.sync.all(x)
* llvm.nvvm.barrier.sync.cnt --> llvm.nvvm.barrier.cta.sync.count(x, y)

Revert "[NVPTX] Unify and extend barrier{.cta} intrinsic support (#140615)"

2025-05-22T17:28:43+00:00

This reverts commit 735209c0688b10a66c24750422b35d8c2ad01bb5.

[NVPTX] Unify and extend barrier{.cta} intrinsic support (#140615)

2025-05-21T15:14:15+00:00

Our current intrinsic support for barrier intrinsics is confusing and
incomplete, with multiple intrinsics mapping to the same instruction and
intrinsic names not clearly conveying intrinsic semantics. Further, we
lack support for some variants. This change unifies the IR
representation to a single consistently named set of intrinsics.

- llvm.nvvm.barrier.cta.sync.aligned.all(i32)
- llvm.nvvm.barrier.cta.sync.aligned(i32, i32)
- llvm.nvvm.barrier.cta.arrive.aligned(i32, i32)
- llvm.nvvm.barrier.cta.sync.all(i32)
- llvm.nvvm.barrier.cta.sync(i32, i32)
- llvm.nvvm.barrier.cta.arrive(i32, i32)

The following Auto-Upgrade rules are used to maintain compatibility with
IR using the legacy intrinsics:

* llvm.nvvm.barrier0 --> llvm.nvvm.barrier.cta.sync.aligned.all(0)
* llvm.nvvm.barrier.n --> llvm.nvvm.barrier.cta.sync.aligned.all(x)
* llvm.nvvm.bar.sync --> llvm.nvvm.barrier.cta.sync.aligned.all(x)
* llvm.nvvm.barrier --> llvm.nvvm.barrier.cta.sync.aligned(x, y)
* llvm.nvvm.barrier.sync --> llvm.nvvm.barrier.cta.sync.all(x)
* llvm.nvvm.barrier.sync.cnt --> llvm.nvvm.barrier.cta.sync(x, y)

[AMDGPU] Set AS8 address width to 48 bits

2025-05-20T00:26:05+00:00

Of the 128-bits of buffer descriptor only 48 bits are address bits, so
following the discussion on https://discourse.llvm.org/t/clarifiying-the-semantics-of-ptrtoint/83987/54,
the logic conclusion is to set the index width to 48 bits instead of
the current value of 128.

Most of the test changes are mechanical datalayout updates, but there
is one actual change: the ptrmask test now uses .i48 instead of .i128
and I had to update SelectionDAGBuilder to correctly extend the mask.

Reviewed By: krzysz00

Pull Request: https://github.com/llvm/llvm-project/pull/139419

[FunctionAttrs] Only consider provenance capture in access attr inference (#138535)

2025-05-06T13:42:26+00:00

For the purpose of inferring readonly/writeonly/readnone on arguments,
we only care about provenance captures, not address captures.

[FunctionAttrs] Add test for SCC with address only capture (NFC)

2025-05-05T13:58:44+00:00

[IR] Relax convergence requirements on call (#135794)

2025-05-02T08:17:58+00:00

Before this commit, having a convergence token on a non-convergent call
was considered to be an error.
This commit relaxes this requirement and allows convergence tokens to be
present on non-convergent calls.

When such token is present, they have no effect as the underlying call
is non-convergent.
This allows passes like DCE to strip `convergent` attribute from
functions for which all convergent operations have been stripped. When
this happens, a convergence token can still exist in the call-site,
causing the verifier to complain.

Alternatives have been considered in #134863 and #134844.

[FunctionAttrs] Bail if initializes range overflows 64-bit signed int (#137053)

2025-04-23T22:56:24+00:00

Otherwise the range doesn't make sense since we interpret it as signed.

Fixes #134115