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)
[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.
[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>
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.
AMDGPU: Migrate GlobalISel CodeGen tests to amdgpu subarch triple (46) (#209754)
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)
[mlir][acc] Introduce acc to gpu codegen pass (#209606)
Introduce pass that lowers `acc.compute_region` to the GPU dialect. For
host-side kernels it wraps the region in `gpu.launch`; for specialized
ACC routines already inside a `gpu.func`, the body is lowered in place.
The pass maps nested `scf.parallel` / `scf.for` loops carrying
`acc.par_dims` to GPU block and thread parallelism, materializes
privatization and reductions for the device, and handles
synchronization.
Example input IR:
```
func.func @two_sibling_loops() {
%c4 = arith.constant 4 : index
%c32 = arith.constant 32 : index
%grid = acc.par_width %c4 {par_dim = #acc.par_dim<block_x>}
%block = acc.par_width %c32 {par_dim = #acc.par_dim<thread_x>}
[94 lines not shown]
AMDGPU: Migrate GlobalISel CodeGen tests to amdgpu subarch triple (45) (#209753)
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 (44) (#209734)
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)
[HIP][SPIR-V] Enable -spirv-preserve-auxdata (#209520)
Function attributes such as `convergent` are only kept through SPIRV if
`-spirv-preserve-auxdata` is enabled. This is important to keep the
translation as accurate as possible. Also, this aligns with the
translator path.
[analyzer] Recognize missing MPI collectives in MPI-Checker (#209263)
Fixes a false positive in the `optin.mpi.MPI-Checker` where legitimate
code was flagged with "Request 'X' has no matching nonblocking call."
Background for reviewers unfamiliar with MPI or this checker
------------------------------------------------------------
MPI (Message Passing Interface) is the standard API for communication
between processes in HPC/cluster programs (C/C++/Fortran). Many MPI
operations come in a *nonblocking* form, prefixed with "I" (for
"immediate"): e.g. `MPI_Isend`, `MPI_Iscatter`. A nonblocking call
returns immediately and hands back an `MPI_Request` handle; the caller
must later complete it with `MPI_Wait` (or `MPI_Waitall`). Forgetting to
wait, or reusing a request before waiting, are real bugs (buffer reuse
races, leaked requests).
The MPI-Checker (originally contributed as a GSoC 2016 project, and is a
domain-specific check - hence it's in "optin").
[57 lines not shown]
[HLSL] Consolidate texture tests (#209348)
Fixes #205878
This PR consolidates all the current texture tests into files that all
texture types can extend by adding new RUN lines.
This should greatly reduce the amount of effort required to implement
and review tests for new texture types.
Assisted by: Claude Opus 4.8
