[CIR] Preserve address spaces in emitPointerWithAlignment casts (#228652)
emitPointerWithAlignment handled CK_AddressSpaceConversion like
CK_BitCast and never emitted the address-space cast, so the result kept
the source address space. When the address escaped (returned reference,
stored pointer), the store path bitcast the destination slot to the
wrong address space. This miscompiled sycl::multi_ptr::operator[] on
SPIR-V: a local-memory offset was used as a generic address. Emit the
cast after the element bitcast, as classic CodeGen does.
createElementBitCast also built the new pointer type in the default
address space, so changing the element type of a non-default-AS pointer
produced an AS-changing bitcast that the verifier rejects (for example
((int *)p)[i] or __builtin_stdc_memreverse8 on an AS3 pointer). Keep the
source address space, which matches classic's withElementType.
Co-authored-by: Claude Opus 5.5 (1M context) <noreply at anthropic.com>
[Clang][Smea] Accept weak reference after declaration
GCC accepts a weak reference even after there is a declaration. It is
fine for us to just append a weak attribute in the previous
declaration directly.
[mlir][XeGPU] Add an SLM round-trip fallback for sg-to-lane convert_layout (#227884)
This PR adds a general fallback lowering for xegpu.convert_layout in the
sg-to-lane distribution pass, which round-trips the value through shared
local memory.
The existing lane-level lowerings are all special cases: the layouts
fold into each other, or they differ in a way a xegpu.lane_shuffle or a
gpu.shuffle can express. Anything else failed to legalize and stopped
the pipeline — for instance a conversion that only moves the lane_data
of a dimension distributed over part of the subgroup:
```mlir
xegpu.convert_layout %src
<{input_layout = #xegpu.layout<lane_layout = [1, 2, 8], lane_data = [1, 1, 4]>,
target_layout = #xegpu.layout<lane_layout = [1, 2, 8], lane_data = [1, 1, 1]>}>
: vector<1x2x32xbf16>
```
[17 lines not shown]
[mlir][xegpu] Fill split-group inner dims on shape_cast result layout (#227560)
This PR set the result layout for shape_cast when it can't infer the
layout from incoming consumer layout. For example,
```mlir
vector.shape_cast %0 : vector<16x1024xbf16> to vector<16x32x32xbf16>
// consumer layout: inst_data = [1, 2, 8], lane_layout = [1, 2, 8], lane_data = [1, 1, 1]
```
Taking the consumer layout as is doesn't work for layout propagation.
The lanes cover 2 along dim 1 and only 8 of dim 2's 32 — a 2x8 box.
Collapsing dims 1 and 2 gives inst_data = [1, 16] on the 16x1024 source,
which is 16 contiguous elements. Those are not the same elements: the
2x8 box is two runs of 8, 32 apart.
This PR fixes the issue by stretching the lane_data to cover the
innermost dims of source shape, to set up result layout so that each
lane gets the same data before and after the shapecast, which is the
condition that shapecast op can be successfully distributed.
```
[10 lines not shown]
[NFC][TSan] Move RestoreStack before ScopedReport in ReportDestroyLocked (#228643)
Run RestoreStack in its own lock scope before constructing ScopedReport
in ReportDestroyLocked (matching ReportRace). This avoids acquiring
ScopedErrorReportLock or symbolizing the current stack if RestoreStack
fails, and avoids holding slot_lock and slot_mtx while populating the
ScopedReport.
Assisted-by: Gemini
T4: boot gates run on the shared nextbsd-ci harness @v0.2.1 (login-only)
img-boot-test.sh + iso-boot-test.sh become thin entry points onto the shared
harness (NB_LOGIN_ONLY; the live ISO adds NB_MEDIA=cd): they extract the
zipped artifact, check out nextbsd/nextbsd-ci at v0.2.1 (pinned tag = no drift),
and run harness/boot-test.sh. The loader un-mute dance, the arch-aware qemu
argv, login detection and teardown now come from the shared harness (one place,
every arch, kept green by its selftest) instead of the in-repo copies.
Delete the 1512-line tests/boot-test.sh monolith (nothing in CI invoked it) and
the in-repo tests/loader.exp.inc + tests/qemu-arch.sh (superseded by the
harness's contract.exp.inc + qemu-arch.sh). The jobs' serial-log dumps point at
the harness transcript. Both gates stay NON-GATING as before (not in release's
needs); the gate is the harness exit class.
Bump the shared-harness pin to v0.2.2 in the thin img/iso boot wrappers
v0.2.2 = v0.2.1 (v0.2.0 + login-only) plus a capped post-banner login-wait
window (180s) that probes for a shell quickly instead of waiting out the whole
global budget on an autologin image (the arm64 30-min 'hang').
[CodeGen] Declare command line options in TableGen
Move the cl::opts of TargetPassConfig.cpp and CodeGenPrepare.cpp into
CodeGenOptions.td, private to lib/CodeGen; other files will follow.
-enable-machine-outliner (cl::ValueOptional) and -regalloc
(RegisterPassParser) stay cl::opt.
CodeGenPrepare and its addressing-mode helpers hold
`const CodeGenOptions &Opts`; TargetPassConfig functions read
CodeGenOptions::Global. getCGPassBuilderOption() converts
std::optional<bool> members to the cl::boolOrDefault fields of the
public CGPassBuilderOption. -basic-block-section-match-infer, which was
not cl::Hidden, is now listed by -help-hidden only.
Aided by Opus 5.5
[X86] Correct the SDTypeProfile for VNNI instructions. (#228745)
Operand 1 and 2 have i8 or i16 elements while the result and operand 0
have i32 elements. The type profile previously said all operands were
the same type.
I've split the type profile to capture the i8 and i16 element size
accurately.
Assisted-by: Claude
[mlir][nvgpu] Align integer WGMMA type checks with i8 operands (#212215)
Integer WGMMA uses i8 operands with an i32 accumulator, but the NVGPU
verifier and K-shape selection still checked for i16.
Changes both checks to i8 and adds tests for rejected i16 inputs and the
existing i8 "not supported yet" limitation.
[NFC][TSan] Move ObtainCurrentStack out of ThreadRegistryLock scope
ObtainCurrentStack only reads the current thread's shadow stack and
allocates a VarSizeStackTrace buffer, which does not require
ThreadRegistryLock (and already runs outside ThreadRegistryLock in
ReportRace, SignalUnsafeCall, and ReportErrnoSpoiling).
Move ObtainCurrentStack (and dummy_pc in ReportDeadlock) before
ScopedReport in ReportMutexHeldWrongContext, ReportMutexMisuse,
ReportDeadlock, and ReportDestroyLocked so the stack trace buffers also
outlive ScopedReport and OutputReport.
Assisted-by: Gemini
Pull Request: https://github.com/llvm/llvm-project/pull/228794
[NFC][TSan] Move RestoreStack before ScopedReport in ReportDestroyLocked
Run RestoreStack in its own lock scope before constructing ScopedReport
in ReportDestroyLocked (matching ReportRace). This avoids acquiring
ScopedErrorReportLock or symbolizing the current stack if RestoreStack
fails, and avoids holding slot_lock and slot_mtx while populating the
ScopedReport.
Assisted-by: Gemini
Pull Request: https://github.com/llvm/llvm-project/pull/228643
[TSan] Defer symbolization in AddStack and AddSleep to SymbolizeStackElems
PR #151495 delayed symbolization for memory accesses, locations,
threads, and mutexes until OutputReport (after ThreadRegistryLock is
released), but missed ScopedReport::AddStack and ScopedReport::AddSleep.
As a result, ReportRace (via AddSleep) and ReportMutexMisuse,
ReportDeadlock, and ReportDestroyLocked (via AddStack) still invoked the
symbolizer while holding ThreadRegistryLock.
Store the unsymbolized stack traces and sleep stack ID in ReportDesc and
symbolize them in ScopedReport::SymbolizeStackElems().
Assisted-by: Gemini
Pull Request: https://github.com/llvm/llvm-project/pull/228795
[NFC][TSan] Allocate ScopedReport as a stack variable
Now that ScopedReport is constructed before acquiring ThreadRegistryLock
or slot locks across all reporting functions, it no longer needs to be
constructed inside the lock scope via placement new on __builtin_alloca
storage.
Declare ScopedReport as a normal stack variable before the lock scope
and remove the manual destructor calls.
Assisted-by: Gemini
Pull Request: https://github.com/llvm/llvm-project/pull/228637
[NFC][TSan] Move ObtainCurrentStack out of ThreadRegistryLock scope
ObtainCurrentStack only reads the current thread's shadow stack and
allocates a VarSizeStackTrace buffer, which does not require
ThreadRegistryLock (and already runs outside ThreadRegistryLock in
ReportRace, SignalUnsafeCall, and ReportErrnoSpoiling).
Move ObtainCurrentStack (and dummy_pc in ReportDeadlock) before
ScopedReport in ReportMutexHeldWrongContext, ReportMutexMisuse,
ReportDeadlock, and ReportDestroyLocked so the stack trace buffers also
outlive ScopedReport and OutputReport.
Assisted-by: Gemini
Pull Request: https://github.com/llvm/llvm-project/pull/228794
[TSan] Lock ScopedErrorReportLock before slot and thread_registry locks
OutputReport runs while ScopedErrorReportLock is held after slot_mtx and
thread_registry have been unlocked. Because code executed during
OutputReport (symbolizer, callbacks, or signal handlers) can acquire
slot_mtx or thread_registry, ScopedErrorReportLock must precede slot and
thread_registry locks in the lock hierarchy to avoid AB-BA deadlocks
between concurrent reports or fork().
- Move ScopedErrorReportLock::Lock() before slot.mtx, thread_registry,
and slot_mtx in ForkBefore (and unlock in reverse order in ForkAfter).
- Replace ctx->thread_registry.CheckLocked() in ScopedReportBase's
constructor with CheckedMutex::CheckNoLocks(), and add CheckLocked() to
AddThread(const ThreadContext *) and CheckNoLocks() to OutputReport.
- Construct ScopedReport before acquiring ThreadRegistryLock across all
reporting functions, and close the RestoreStack lock scope before
constructing ScopedReport in ReportRace.
Assisted-by: Gemini
[2 lines not shown]
[NFC][TSan] Move RestoreStack before ScopedReport in ReportDestroyLocked
Run RestoreStack in its own lock scope before constructing ScopedReport
in ReportDestroyLocked (matching ReportRace). This avoids acquiring
ScopedErrorReportLock or symbolizing the current stack if RestoreStack
fails, and avoids holding slot_lock and slot_mtx while populating the
ScopedReport.
Assisted-by: Gemini
Pull Request: https://github.com/llvm/llvm-project/pull/228643
[NFC][TSan] Use in-class member initializers in tsan_report.h
Use in-class member initializers for all structs and classes in
tsan_report.h and default constructors and destructor in
tsan_report.cpp.
Assisted-by: Gemini
Pull Request: https://github.com/llvm/llvm-project/pull/228780
[orc-rt] Make noErrors report errors as test failures (#228807)
noErrors used cantFail, so an unexpected error aborted the whole test
binary without naming the test that hit it, and with assertions disabled
was dropped silently. Use EXPECT_THAT_ERROR instead, so the error is
reported as a failure of the current test and the remaining tests still
run.
rpc_generic.c: Initialize "cp" to shut the compiler up
This patch does not fix any semantics issue.
MFC after: 3 months
Fixes: 884ee8d6c9b4 ("nfscl: Add some glue for client side NFS over RDMA")
[NFC][TSan] Merge ScopedReportBase into ScopedReport (#228640)
ScopedReport is the only subclass of ScopedReportBase and nothing uses
ScopedReportBase directly. Merge ScopedReportBase into ScopedReport.
Assisted-by: Gemini
[SLP]Fix crash on deleted main op of copyable state in reductions
Vectorizing one group of reduced values can erase the main operation of
a later group's copyable state, leaving it with dropped operands.
Fixes #228768
Reviewers:
Pull Request: https://github.com/llvm/llvm-project/pull/228806
[orc-rt] Use the Error matchers in SPSWrapperFunctionTest (#228804)
Use the Error matchers introduced in 4c8a437d0487 to clean up error
tests in SPSWrapperFunctionTest.
Re-enable Rescan when a WiFi scan fails
A failure in networkdictionary() inside the rescan thread, or the
WiFi card vanishing from the new scan, left the Rescan button greyed
out for good. Restore it in a finally block and skip the access point
list update when the card is gone.
[NFC][TSan] Move ObtainCurrentStack out of ThreadRegistryLock scope
ObtainCurrentStack only reads the current thread's shadow stack and
allocates a VarSizeStackTrace buffer, which does not require
ThreadRegistryLock (and already runs outside ThreadRegistryLock in
ReportRace, SignalUnsafeCall, ReportErrnoSpoiling, and
ReportDestroyLocked).
Move ObtainCurrentStack (and dummy_pc in ReportDeadlock) before
ScopedReport in ReportMutexHeldWrongContext, ReportMutexMisuse, and
ReportDeadlock so the stack trace buffers also outlive ScopedReport and
OutputReport.
Assisted-by: Gemini
Pull Request: https://github.com/llvm/llvm-project/pull/228794
[NFC][TSan] Move RestoreStack before ScopedReport in ReportDestroyLocked
Run RestoreStack in its own lock scope before constructing ScopedReport
in ReportDestroyLocked (matching ReportRace). This avoids acquiring
ScopedErrorReportLock or symbolizing the current stack if RestoreStack
fails, and avoids holding slot_lock and slot_mtx while populating the
ScopedReport.
Assisted-by: Gemini
Pull Request: https://github.com/llvm/llvm-project/pull/228643
