BSD Commit Log Search

AMDGPU: Migrate GlobalISel CodeGen tests to amdgpu subarch triple (53) (#209779)

Fold the explicit -mcpu subtarget into the amdgpu subarch triple on
llc/opt RUN lines (e.g. -mtriple=amdgcn -mcpu=gfx600 ->
-mtriple=amdgpu6.00), now that these tests carry an explicit -mcpu
following the default-subtarget pinning. Codegen output is unchanged.

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

[LV] Fix LoopVectorize/X86/CostModel tests when regenerated with UTC. NFC (#209802)

Some of the cost model tests were passing, but their UTC arguments
were outdated and CHECK lines disappeared when regenerating them.

- Add additional --filter lines for VPlan cost model CHECK lines,
  which were manually written previously
- Update the --filter argument for the updated INTERLEAVE-GROUP print
  output
- Regenerate strided-load-*.ll tests with UTC
- Pass -vplan-print-metadata=false for tests that were using
  hand-written regexes in CHECK lines

[gn build] Port ef118b5b3a61 (#209800)

AMDGPU: Migrate CodeGen tests to amdgpu subarch triple (52) (#209778)

Fold the explicit -mcpu subtarget into the amdgpu subarch triple on
llc/opt RUN lines (e.g. -mtriple=amdgcn -mcpu=gfx600 ->
-mtriple=amdgpu6.00), now that these tests carry an explicit -mcpu
following the default-subtarget pinning. Codegen output is unchanged.

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

AMDGPU: Migrate CodeGen tests to amdgpu subarch triple (51) (#209777)

Fold the explicit -mcpu subtarget into the amdgpu subarch triple on
llc/opt RUN lines (e.g. -mtriple=amdgcn -mcpu=gfx600 ->
-mtriple=amdgpu6.00), now that these tests carry an explicit -mcpu
following the default-subtarget pinning. Codegen output is unchanged.

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

[gn build] Port 9a4deccf90ae (#209799)

[gn build] Port 2375505c7c19 (#209797)

[clang][bytecode] Fix virtual-bases.cpp on 32 bit hosts (#209766)

It currently fails on clang-armv8-quick with:
```console
# | error: diagnostics with 'error' severity seen but not expected: 
# |   File /home/tcwg-buildbot/worker/clang-armv8-quick/llvm/clang/test/AST/ByteCode/virtual-bases.cpp Line 441: static assertion failed due to requirement '((__builtin_constant_p((char *)&c.c) ? ((char *)&c.c) : ((char *)&c.c)) - (__builtin_constant_p((char *)&c) ? ((char *)&c) : ((char *)&c))) == 12'
# |   File /home/tcwg-buildbot/worker/clang-armv8-quick/llvm/clang/test/AST/ByteCode/virtual-bases.cpp Line 442: static assertion failed due to requirement '((__builtin_constant_p((char *)&c.b) ? ((char *)&c.b) : ((char *)&c.b)) - (__builtin_constant_p((char *)&c) ? ((char *)&c) : ((char *)&c))) == 8'
# |   File /home/tcwg-buildbot/worker/clang-armv8-quick/llvm/clang/test/AST/ByteCode/virtual-bases.cpp Line 443: static assertion failed due to requirement '((__builtin_constant_p((char *)&c.a) ? ((char *)&c.a) : ((char *)&c.a)) - (__builtin_constant_p((char *)&c) ? ((char *)&c) : ((char *)&c))) == 20'
# |   File /home/tcwg-buildbot/worker/clang-armv8-quick/llvm/clang/test/AST/ByteCode/virtual-bases.cpp Line 444: static assertion failed due to requirement '((__builtin_constant_p((char *)&c.x) ? ((char *)&c.x) : ((char *)&c.x)) - (__builtin_constant_p((char *)&c) ? ((char *)&c) : ((char *)&c))) == 16'
# | error: diagnostics with 'note' severity seen but not expected: 
# |   File /home/tcwg-buildbot/worker/clang-armv8-quick/llvm/clang/test/AST/ByteCode/virtual-bases.cpp Line 441: expression evaluates to '8 == 12'
# |   File /home/tcwg-buildbot/worker/clang-armv8-quick/llvm/clang/test/AST/ByteCode/virtual-bases.cpp Line 442: expression evaluates to '4 == 8'
# |   File /home/tcwg-buildbot/worker/clang-armv8-quick/llvm/clang/test/AST/ByteCode/virtual-bases.cpp Line 443: expression evaluates to '16 == 20'
# |   File /home/tcwg-buildbot/worker/clang-armv8-quick/llvm/clang/test/AST/ByteCode/virtual-bases.cpp Line 444: expression evaluates to '12 == 16'
# | 8 errors generated.
```

AMDGPU: Migrate CodeGen tests to amdgpu subarch triple (50) (#209776)

Fold the explicit -mcpu subtarget into the amdgpu subarch triple on
llc/opt RUN lines (e.g. -mtriple=amdgcn -mcpu=gfx600 ->
-mtriple=amdgpu6.00), now that these tests carry an explicit -mcpu
following the default-subtarget pinning. Codegen output is unchanged.

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

AMDGPU: Migrate GlobalISel CodeGen tests to amdgpu subarch triple (49) (#209757)

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)

Revert "[SLP] Support memory runtime alias checks"

This reverts commit 96a53b1c5356754b4708c0fee8517fd0d69cb1fd to fix
buildbots https://lab.llvm.org/buildbot/#/builders/223/builds/7209
reported in https://github.com/llvm/llvm-project/pull/209305?email_source=notifications&email_token=ABI45DXMMB4TIQHXXWA6NYD5E57EVA5CNFSNUABFM5UWIORPF5TWS5BNNB2WEL2JONZXKZKDN5WW2ZLOOQXTIOJYGA4TIMZVGU3KM4TFMFZW63VMON2GC5DFL5RWQYLOM5S2KZLWMVXHJLDGN5XXIZLSL5RWY2LDNM#issuecomment-4980943556

Reviewers: 

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

Rebase

Created using spr 1.3.7

[AMDGPU][CodeGen] Place `S_NOP` after `S_SETREG_IMM32_B32` in predecessor MBB (#209620)

When placing a `S_SET_VGPR_MSB` at the beginning of a block, check
whether there exists a `S_SETREG_IMM32_B32(MODE)` in the predecessor
block that falls-through into it; if there is one a `S_NOP` has to be
inserted in between them.

[compiler-rt] Fix amdgcn -> amdgpu conversion for compiler-rt (#209765)

Summary:
We are apparently transitioning to using `amdgpu` everywhere, so these
need to be updated.

[Clang][AMDGPU] Search both amdgcn and amdgpu libraries (#209770)

Summary:
Temporary workaround as we transition triples from amdgcn to amdgpu.

[AMDGPU] Avoid errors with 24-bit div/rem (#209711)

Avoid errors when using floating-point reciprocal to calculate Y/X when
Y = (0x7FFFFF/X)*X-1. Limit expansion to 23-bit signed and 22-bit
unsigned to avoid this issue.

This change is analogous to the change done in
https://github.com/llvm/llvm-project/pull/202753, but in
AMDGPUISelLowering.cpp.

Signed-off-by: John Lu <John.Lu at amd.com>

[lldb][bazel] Add ObjectFileXCOFF plugin library to the Bazel overlay (#209747)

Adds the PluginObjectFileXCOFF cc_library to the lldb Bazel overlay,
modeled on PluginObjectFileELF/PluginObjectFileCOFF and mirroring
source/Plugins/ObjectFile/XCOFF/CMakeLists.txt (lldbCore, lldbHost,
lldbSymbol, lldbTarget; llvm BinaryFormat/Object/Support). The library
is intentionally left out of DEFAULT_PLUGINS.

bazel rule assisted with: claude

Can confirm this library will convert to BUCK internally at Meta and
build with buck2 where we do have in the default plugin list. I left it
out here, unless there's a desire to add it in.

[lldb][bazel] Add PlatformWasm plugin library to the Bazel overlay (#209749)

Adds the PluginPlatformWasm cc_library plus its PlatformWasmProperties
gentbl_cc_library (settings .td, modeled on PlatformQemuUserProperties)
to the lldb Bazel overlay. Deps mirror
source/Plugins/Platform/WebAssembly/CMakeLists.txt (lldbCore, lldbHost,
lldbTarget, lldbUtility, Support) plus the :PluginPlatformGDB and
:PluginProcessWasm plugin deps its GDB-remote/Web-Inspector sources use.
The library is intentionally left out of DEFAULT_PLUGINS.

bazel rule assisted with: claude

Can confirm this library will convert to BUCK internally at Meta and
build with buck2 where we do have in the default plugin list. I left it
out here, unless there's a desire to add it in.

[CIR][AMDGPU] Add support for AMDGCN sqrt builtins (#197342)

Adds codegen for the following AMDGCN square root builtins:

- __builtin_amdgcn_sqrt (double)
- __builtin_amdgcn_sqrtf (float)
- __builtin_amdgcn_sqrth (half)
- __builtin_amdgcn_sqrt_bf16 (bfloat16)

These are lowered to the corresponding `llvm.amdgcn.sqrt` intrinsic calls.

[mlir][acc] Add attribute for redundant execution marking (#209784)

Adds `acc.gpu_block_redundant` attribute for correct gang-redundant
lowering.

OpenACC gang-redundant regions execute on every gang rather than
partitioning work across gangs. When lowering to the GPU dialect, such
regions must not be assigned block-level par dims; only thread-level
parallelism applies. This marks loops with this attribute to ensure
proper execution and to avoid incorrect predication.

Co-authored-by: Delaram Talaashrafi <dtalaashrafi at nvidia.com>

[NFC][OpenACC] fix doc formatting for reduction ops (#209739)

Was not printing correctly due to the matching `$`

AMDGPU: Migrate CodeGen tests to amdgpu subarch triple (52)

Fold the explicit -mcpu subtarget into the amdgpu subarch triple on llc/opt RUN lines (e.g. -mtriple=amdgcn -mcpu=gfx600 -> -mtriple=amdgpu6.00), now that these tests carry an explicit -mcpu following the default-subtarget pinning. Codegen output is unchanged.

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

AMDGPU: Migrate tests with regenerated checks to amdgpu subarch triple (54)

Fold the explicit -mcpu subtarget into the amdgpu subarch triple on autogenerated tests where the folded triple changes output relative to the previous default subtarget (e.g. cost-model BASE lines, scheduling). CHECK lines were regenerated with the update_*_test_checks.py scripts.

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

AMDGPU: Migrate GlobalISel CodeGen tests to amdgpu subarch triple (53)

Fold the explicit -mcpu subtarget into the amdgpu subarch triple on llc/opt RUN lines (e.g. -mtriple=amdgcn -mcpu=gfx600 -> -mtriple=amdgpu6.00), now that these tests carry an explicit -mcpu following the default-subtarget pinning. Codegen output is unchanged.

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

AMDGPU: Migrate CodeGen tests to amdgpu subarch triple (51)

Fold the explicit -mcpu subtarget into the amdgpu subarch triple on llc/opt RUN lines (e.g. -mtriple=amdgcn -mcpu=gfx600 -> -mtriple=amdgpu6.00), now that these tests carry an explicit -mcpu following the default-subtarget pinning. Codegen output is unchanged.

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

AMDGPU: Migrate CodeGen tests to amdgpu subarch triple (50)

Fold the explicit -mcpu subtarget into the amdgpu subarch triple on llc/opt RUN lines (e.g. -mtriple=amdgcn -mcpu=gfx600 -> -mtriple=amdgpu6.00), now that these tests carry an explicit -mcpu following the default-subtarget pinning. Codegen output is unchanged.

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

AMDGPU: Migrate GlobalISel CodeGen tests to amdgpu subarch triple (49)

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)

AMDGPU: Migrate GlobalISel CodeGen tests to amdgpu subarch triple (48) (#209756)

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)

AMDGPU: Migrate GlobalISel CodeGen tests to amdgpu subarch triple (47) (#209755)

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)

[Flang][OpenMP] Implement DEPEND clause for TASKWAIT directive (#193568)

Implements support for the `DEPEND` clause on `TASKWAIT`, added in
OpenMP 5.0.