BSD Commit Log Search

[AMDGPU] Model structured control-flow lane masks as i1

AMDGPU: Migrate CodeGen tests to amdgpu subarch triple (14) (#209117)

Mechanically migrate the command-line target spelling on llc/opt RUN
lines in
llvm/test/CodeGen/AMDGPU from -mtriple=amdgcn ... -mcpu=<gfx> to the
folded
subarch triple form (e.g. -mtriple=amdgpu9.00-amd-amdhsa), dropping the
redundant -mcpu.

Co-Authored-By: Claude <noreply at anthropic.com> (Claude-Opus-4.8)

Update flag.

[AArch64][SVE] Support lowering masked loads of <4 x bf16> and <8 x bf16>

Add support for lowering masked loads of <4 x bf16> and <8 x bf16> when
target features contain "+sve".

AMDGPU: Migrate CodeGen tests to amdgpu subarch triple (13) (#208897)

Mechanically migrate the command-line target spelling on llc/opt RUN
lines in
llvm/test/CodeGen/AMDGPU from -mtriple=amdgcn ... -mcpu=<gfx> to the
folded
subarch triple form (e.g. -mtriple=amdgpu9.00-amd-amdhsa), dropping the
redundant -mcpu.

Co-Authored-By: Claude <noreply at anthropic.com> (Claude-Opus-4.8)

[review] Improve test.

[lldb] Correct alignment of register fields (#209140)

In #188049 the alignment of register names for the raw value part "pc =
0x..." changed.

This resulted in many register fields being 2 spaces off where they
should be, because the stream we were printing to had a 2 space indent
set where it had none before.

This is actually a bug in the original field formatting code. It assumed
the stream had no indent level of its own. So the fix is to indent using
the stream (the 2 spaces), then indent to where the register name ends.

[reviews] Improve test.

fixup! Okay, change PR so we just remover the gating from sys reg and leave FEAT_HCX in the code

[AMDGPU] Add pre-commit test demonstrating current control-flow lowering (NFC)

[AArch64][llvm][tablegen] Restrict luti6 assembly (4 regs, 8-bit) to 0 <= Zn <= 7 (#200751)

The `luti6` instruction (four registers, 8-bit) should only allow
assembly of `0 <= Zn <= 7`, since there's only 3 bits for `Zn`. It
actually allows > 7:
```
   luti6 { z0.b - z3.b }, zt0, { z8 - z10 }
```
which produces a duplicate encoding to the following:
```
   luti6 { z0.b - z3.b }, zt0, { z0 - z2 }
```

Update tablegen to handle the inferred register-class naming
collision caused by adding the explicit z0-z7 ZPR3 class.

[RISCV] Support spacemit vsmtvdotii extensions (#202533)

SPEC: https://github.com/spacemit-com/docs-ai/blob/main/en/architecture/ime_extension.md

[AMDGPU] Remove option to disable SGPR hazard waits (#208720)

This was useful when initially implemented, but probably be removed now.

clang-linker-wrapper: Use AMDGPU::TargetID for device-image compatibility

Previously the link compatibilty was based on a raw string match of the
triple + cpu name. Start parsing the triple, and use the compatibility
function. Start using AMDGPU::TargetID from TargetParser, as a step towards
removing clang's redundant TargetID handling.

Co-authored-by: Claude (Opus 4.8)

compiler-rt: ubsan: Add suppressions test for nested functions (#206962)

This is a test-only change to increase the test coverage of UBSan to
clarify that named suppressions only apply to the function itself and do
not cover (possibly inlined) called functions.

The background is that, while trying to fix
https://github.com/llvm/llvm-project/issues/132533, I discovered missing
test coverage and possibly confused myself into a wrong fix in
https://github.com/llvm/llvm-project/pull/206735#discussion_r3505816693

So I think it could make sense to add test coverage for the case that I
broke, which serves as a reference baseline ground truth.

---------

Co-authored-by: MarcoFalke <*~=`'#}+{/-|&$^_ at 721217.xyz>

[review] improve test.

[Polly] Fix infinite recursion in `ScopExpander::visitAddRecExpr` (#206063)

visitUnknown follow VMap, call GenSE.getSCEV() on the mapped value, and
get back the original AddRec, recursing visitAddRecExpr infinitely.

Insert the result into SCEVCache before calling visit() to cut
recursion.

Fixes #205137.

[VPlan] Expand SCEV min/max via scalar intrinsic VPInstructions. (#207842)

Use the VPInstruction::Intrinsic opcode to expand min/max SCEV
expressions to LLVM min/max intrinsics as VPInstructions.

The operands are chained in reverse order, matching SCEVExpander's
expansion of min/max expressions.

This helps to clean up another set of duplicated vscale calls.

Depends on https://github.com/llvm/llvm-project/pull/207541
Depends on https://github.com/llvm/llvm-project/pull/207836

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

[RISCV] Add TuneNoDefaultUnroll to spacemit-x100 (#209069)

In the spec2k6int tests, some subtests improved while others
regressed. The overall score shows a 0.7% gain. Given that
most OoO CPUs enable this feature by default, also enable
it by default for X100.

[reivews] Simplify processing.

[CIR][AMDGPU] Add support for AMDGCN rcp builtins (#197447)

Adds codegen for the following AMDGCN reciprocal builtins:

- __builtin_amdgcn_rcp (double)
- __builtin_amdgcn_rcpf (float)
- __builtin_amdgcn_rcph (half)
- __builtin_amdgcn_rcp_bf16 (bfloat16)

These are lowered to the corresponding `llvm.amdgcn.rcp` intrinsic.

Add other missing include

Re-add include

[CodeGen][AMDGPU] Prepare rematerializer for subreg remat support (NFC)

This makes some NFCs to the rematerializer before starting to improve
support for sub-register rematerialization. The main changes are the
replacement of `Rematerializer::Reg::Dependency` type (essentially a
pair of a machine operand index and a register index) in favor of a
simple register index, dropping the machine operand index. The latter
has no current uses and will lose meaning once we allow rematerializable
registers to be defined by multiple MIs. Similarly, and for the same
rationale, unrematerializable register dependencies are now tracked as
a register/lanemask pair instead of a machine operand index.

Other minor changes listed below.

- Removal of `DefRegion` argument to `Rematerializer::recreteReg`.
  Registers are always re-created in their original region so there is
  no need to set their region again.
- Removal of `InsertPos` unused argument to
  `Rematerializer::postRematerialization`.

    [3 lines not shown]

[clang][bytecode] Improve diagnostics on typeid field access (#209101)

Do some extra work to extract the requested field name.

Use a union instead of void* + static_cast for InsertBeforePos

[offload][test] Add fallback to legacy AMDGPU triple library dir

After the amdgcn-amd-amdhsa -> amdgpu-amd-amdhsa rename, look up the
device runtime libraries in the legacy directory when the runtime was
built under the old triple name.

Claude assisted with this patch.

[AMDGPU] Use v_pk_add_bf16 for scalar bf16 fadd on gfx1250/gfx13

[NFC][SPIR-V] Remove unused LowLevelType parameter from buildLoadInst (#208428)

[mlir][SPIR-V] Add SPIRVToLLVM conversion for SNegate (#206950)