[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.
[AMDGPU] Match bitsin(typeof(x)) - ctpop(x) to s_bcnt0_i32 (#164847)
This PR optimizes the pattern bitsin(typeof(x)) - popcnt(x) to
s_bcnt0_i32 on AMDGPU. It also creates a Clang builtin for s_bcnt0_i32
so that users can call this instruction directly instead of relying on
the compiler to match this pattern.
[offload][lit] Disable failing olCreateProgram test on AMD (#213126)
Fails with an error about not being able to find `lld`. AMD team is
investigating, but disable the test in the meantime.
Context: https://github.com/llvm/llvm-project/pull/212860
Signed-off-by: Nick Sarnie <nick.sarnie at intel.com>
X86: Simplify the EH_LABEL Expand condition
Re-express the opt-out handling of EH_LABEL. The special
case is 32-bit non-GNU Windows, and the net result is to skip
printing unused labels. Try to make this more comprehensible
to help figure out where this logic should really be.
I want to eliminate use of the TargetOptions::ExceptionModel,
which ideally wouldn't be needed in a TargetLowering
constructor.
Co-authored-by: Claude (Claude-Opus-4.8) <noreply at anthropic.com>
