AMDGPU: Give v_cvt_sr_pk_bf16_f32 its own subtarget feature
v_cvt_sr_pk_bf16_f32 was gated on bf16-cvt-insts, but that feature is
also present on gfx950 where the (non-sr) v_cvt_pk_bf16_f32 was first
added. The stochastic-rounding v_cvt_sr_pk_bf16_f32 was only added
for gfx1250 and has no gfx950 encoding, so it would mis-select and
later hit the "Invalid opcode" assert. Introduce cvt-sr-pk-bf16-f32-inst,
currently added to gfx13 and 125*
Co-authored-by: Claude (Claude-Opus-4.8)
[clang-format] Prevent re-assigning type on finalized tokens (#210763)
Prevents a finalized token inside modifyContext from being reassigned
through the setType member function by checking if the token is
finalized.
This ensures ill-defined code like does not trigger an assertion failure
during reformatting.
Fixes #210509
[Offload][Test] Add llvm bin in the PATH for offload-unit suite (#213149)
This PR makes the offload-unit suite put llvm bin directory on PATH so
that tests need lld can find. It fixes the issue exposed in:
https://github.com/llvm/llvm-project/pull/212860
[SBVec] Add top-down vectorization to the unified Sandbox Vectorizer
Extend the Sandbox Vectorizer's `bottom-up-vec` pass so a single
implementation can vectorize in either direction, and add the top-down
strategy that walks def-use chains forward from a seed.
Direction selection
--------------------
The pass direction is chosen from the Region's auxiliary pass argument:
"bottom-up" (or empty, the default) and "top-down" map onto a
SchedDirection, and any other value is rejected with a fatal usage error.
The vectorizer always runs in the same direction as the scheduler.
Top-down traversal
------------------
Bottom-up starts from a seed slice (e.g. stores to consecutive addresses)
and recurses into operands. Top-down instead starts from a seed of
consecutive loads and recurses into *users*:
[32 lines not shown]
[SandboxIR] Fix notifyEraseInstr to skip scheduled neighbors
Guard both loops with !PredN->scheduled() / !SuccN->scheduled() so
scheduled neighbors are left untouched, and add a unit test that erases
a node with one scheduled and one unscheduled predecessor to cover the
fix.
[AMDGPU] Use a single SubtargetPredicate for fp8/bf8 -> f32 conversions (#212888)
Add FeatureCvtFP8SDWASrcSel for the form that takes the byte from the SDWA src0_sel field and FeatureCvtFP8ByteSel for the form that takes it from a byte_sel operand. Both imply FeatureFP8ConversionInsts, and a subtarget provides one of them, so the multiclass-generated HasCvtFP8SDWASrcSel and HasCvtFP8ByteSel replace the GFX9 and GFX11Plus predicates outright.
Assisted-By: Claude Opus 5
[flang] Add a pass to get OpenACC device ptr for CUDA kernel (#212299)
When a CUDA kernel is launched inside an OpenACC data region, it does
not properly get the device pointers and instead uses host data. This PR
adds a pass to get the device pointers set up by the OpenACC data
construct and pass them explicitly to the CUDA kernel.
Note: A possibly non-contiguous array argument is currently not supported. This pass will skip these cases.
---------
Co-authored-by: Yebin Chon <ychon at nvidia.com>
[mlir][OpenACC] Atomicize contended shared array reduction updates (#212971)
Example:
```fortran
!$acc parallel loop gang reduction(+:a)
do i = 1, N
!$acc loop worker reduction(+:a)
do j = 1, M
a(i) = a(i) + b(j,i)
end do
end do
```
In this code, the worker accumulator is block-shared, so every worker
updates
the same element with a plain read-modify-write and all but one partial
is lost.
Fix: make in-place updates of a block-shared array accumulator atomic,
[7 lines not shown]
[LV] Add tests for preserving branch weights through VPlan (NFC). (#213131)
Add coverage for carrying the branch weights of a predicated block from
VPlan0 through to the generated IR
[clang][DependencyScanning] Extracting a Driver-free CompilerInstaneWithContext Initializer (#211405)
This PR extracts out an initializer of `CompilerInstanceWithContext`
that does not depend on any driver code in preperation of moving the
`CompilerInstanceWithContext` into `DependencyScanningWorker.cpp` as an
implementation detail that is not exposed by any APIs.
---
<sub>Stack created with <a
href="https://github.com/github/gh-stack">GitHub Stacks CLI</a> • <a
href="https://gh.io/stacks-feedback">Give Feedback 💬</a></sub>
AMDGPU: Add missing msad-insts to gfx13 frontend feature map
fillAMDGCNFeatureMap omitted msad-insts for gfx1310/gfx13-generic, so
clang wrongly rejected __builtin_amdgcn_msad_u8 on those targets even
though the backend enables the feature (FeatureGFX9 generation, inherited
through FeatureGFX13). Add it to the gfx13 case.
Co-authored-by: Claude (Claude-Opus-4.8)
AMDGPU: Add missing flat-global-insts to gfx13 frontend feature map
fillAMDGCNFeatureMap omitted flat-global-insts for gfx1310/gfx13-generic,
so clang wrongly rejected __builtin_amdgcn_av_load_b128 /
__builtin_amdgcn_av_store_b128 on those targets even though the backend
enables the feature. Add it to the gfx13 case.
Co-authored-by: Claude (Claude-Opus-4.8)
[clang-format] Add SpacesInComments option for block comments (#204727)
Adds a new `SpacesInBlockComments` clang-format option to control
spacing after `/*` and before `*/` in ordinary block comments.
Supported values:
* Always: formats `/*comment*/` as `/* comment */`
* Never: formats `/* comment */` as `/*comment*/`
* Leave: preserves existing spacing
Documentation comments such as `/** ... */` and `/*! ... */`, and
parameter comments such as `/*Arg=*/`, are left unchanged.
Tests added for all option values, multiline block comments, and
excluded parameter/doc comments.
Addresses #160682
[lldb/script] Migrate synthetic children providers onto ScriptedPythonInterface (#210845)
Give `type synthetic add -l` a formal
`ScriptedSyntheticChildrenInterface`, matching the architecture used
elsewhere in this series: a C++ interface header, a Python-backed
implementation, `PluginManager` registration with CLI/API usages, and a
generatable ABC template (`scripted_synthetic_children.py`) wired into
`scripting extension generate`.
Every method goes through the shared `Dispatch<T>()` machinery instead
of hand-rolling its own Locker/raw-SWIG calls. `Dispatch<T>()` is taught
to introspect the target method's arity via
`PythonCallable::GetArgInfo()` and drop trailing args before calling, so
providers that legitimately define an argument as optional
(`num_children(self)` vs. `num_children(self, max_count)`) still work
through the generic dispatch path.
This retires the ad-hoc `LLDBSwigPython_*` synthetic-children bridge
functions entirely.
Signed-off-by: Med Ismail Bennani <ismail at bennani.ma>
[lldb] Add a unit test for Wasm architecture compatibility (#213070)
A Wasm module encodes neither a vendor nor an OS, so ObjectFileWasm
reports a bare wasm32 or wasm64 architecture. That keeps an on-disk
module compatible with the more specific triple a Wasm runtime reports
at launch, which lets the dynamic loader reuse the module instead of
reparsing it from process memory. An over-specified triple, including
one that spells out an unknown vendor and OS, does not match, because an
explicitly specified component still counts as specified.
