Mips: Route ABI queries through the subtarget
MipsSubtarget::getABI() already forwards to the target machine, so every
direct MipsTargetMachine::getABI() reader can use the subtarget instead.
Co-Authored-By: Claude Opus 5 <noreply at anthropic.com>
[BPI] Add getEdgeProbabilitiesFromWeights (NFC). (#226560)
Move computing edge probabilities from branch weights to a static
helper. This allows re-using the core logic when computing execution
frequencies in VPlan, fixing a number of subtle differences.
While touching the code, also modernize the code slightly.
Note that the helper does not handle the case where all successors are
unreachable. This would need additional info and is not needed for the
VPlan case.
PR: https://github.com/llvm/llvm-project/pull/226560
[clang-repl] Implement IncrementalHIPDeviceParser for HIP device compilation (#218337)
This PR implements the IncrementalHIPDeviceParser. It takes each
incremental statement's device code and lowers it all the way to a
loadable HIP fat binary: linking the ROCm device libraries, running the
AMDGPU optimization/codegen pipeline, producing a .hsaco code object,
and embedding it as __hip_fatbin for the host runtime to register and
launch.
This PR only adds the parser class, wiring it into Interpreter.cpp and
the end-to-end tests will come in a follow-up.
Assisted by Claude Opus 4.8
[orc-rt] make bind_front match std::bind_front semantics. (#226646)
orc_rt::bind_front stands in for std::bind_front until the ORC runtime
can use C++20, so code written against it should keep working after the
switch. This commit brings its behavior in line with std::bind_front:
- noexcept propagation
- const and rvalue call operators, forwarding the wrapper's qualifiers
- no fallback to a differently-qualified call operator
- reference return types
- member pointers
- copyable wrappers, including nested binds
The unit tests are reorganized and extended to cover these.
Assisted-by: Claude
[LV] Always try to use fixed VF for low TC loops if no epilogue allowed. (#226318)
Remove the MaxPowerOf2RuntimeVF gate for falling back to using fixed
width VFs when no epilogue is allowed and the trip count is below the
minimum for tail folding.
MaxPowerOf2RuntimeVF being not set means we cannot compute the maximum
runtime VF, due to missing max vscale. In that case we are not able to
determine if a epilogue loop remains for scalable VFs, but we can still
pick a fixed VF, if no epilogue remains for it.
Also updates processLoop to ignore CM_EpilogueNotNeededFoldTail, if the
trip count is below the tail-folding threshold.
PR: https://github.com/llvm/llvm-project/pull/226318
Fix for a clang-repl crash when passing non-existent file to --Xcc (#225861)
Problem: When clang-repl fails to initialize because an --Xcc file does
not exist, Interpreter::create() destroys the partially initialized
Interpreter. Its destructor unconditionally calls FinalizeAction(),
which invokes FrontendAction::EndSourceFile() without a registered
CompilerInstance and triggers an assertion.
Fix: Track whether interpreter initialization completed successfully and
only finalize the frontend action for successfully initialized
interpreters.
Test: Add a regression test covering a nonexistent --Xcc file.
This change resolves https://github.com/llvm/llvm-project/issues/225429
Assisted by gpt-5.6-luna
---------
Co-authored-by: Anutosh Bhat <andersonbhat491 at gmail.com>
Revert "[clang][bytecode] Stop relying on `CheckEvaluationResult()` (… (#226637)
…#186045)"
This reverts commit 55b498466a7b2620d09fcdf1b2300f71a35f6d8b.
This breaks two msan builders and reproducing the failure locally takes
forever.
[CIR] Propagate the record address space to get_member
A member is stored inside its record, so a pointer to it is in the same
address space as the record. Several get_member builders created the
result pointer in the default address space regardless of the base:
cir.get_member %p[0] : !cir.ptr<!rec_S, target_address_space(4)>
-> !cir.ptr<!s32i>
On SPIR-V the default address space is private, so a SYCL kernel reading
a captured pointer from its closure (in the generic address space) went
through a private pointer. For unions, lowering emitted a bitcast that
changed the address space, which is invalid LLVM IR.
Take the address space from the base in emitAddrOfFieldStorage,
getAddrOfBitFieldStorage, the Address-based createGetMember, multi-output
inline asm (whose temporary is an AS 5 alloca on AMDGPU) and the NVPTX
printf argument buffer. Add a verifier check so a mismatch fails early.
[CIR] Cast global addresses to their declared address space
A global can live in a different address space than its declared type,
e.g. a CUDA __device__ or __shared__ variable. Classic CodeGen casts the
address once where it is formed (GetOrCreateLLVMGlobal and
getOrCreateStaticVarDecl), so every user sees a pointer in the declared
(generic) address space. CIR kept the global's address space on the value
and relied on each user to cast it. Users that did not, such as returning
or storing the address, bitcast the destination slot instead,
reinterpreting the pointer:
__shared__ int sh;
__device__ int *f() { return &sh; }
stored the raw shared-window address as a generic pointer, with no
cvta.shared on NVPTX. After #226455 the same applies to function-local
__shared__ variables.
Do the same check as classic CodeGen where CIR first has the address as a
[5 lines not shown]
Revert "[clang][OpenMP] Only accept char strings in the message clause" (#226636)
Reverts llvm/llvm-project#226080
The newly added test error_message.c is failing on the PS4 bot
https://lab.llvm.org/buildbot/#/builders/144/builds/66294
[libc++][NFC] Avoid empty namespace in internal headers of `<ranges>` (#226034)
...in pre-C++20/23 modes. This follows up
0c94a978df783bd67535350eac84088fc482716a and fixes more complaints from
clang-tidy checks in CI.
[MLIR][Python] Allow `Attribute` to denote unconstrained attributes in Python-defined ops (#224679)
Operation definitions support attribute constraints such as
`IntegerAttr`, `StringAttr`, and `IntegerAttr | StringAttr`, but
currently reject `Attribute` itself.
This change allows `Attribute` to denote an unconstrained attribute,
accepting both built-in and Python-defined attributes:
```python
class MyOp(MyDialect.Operation, name="my_op"):
arg: Operand
attr: Attribute
res: Result
```
This complements the unconstrained `Operand` and `Result` annotations
introduced in #220216. The annotation also works with `attribute()`
field specifiers, including `kw_only` and `default_factory`.
Assisted-by: Codex / GPT6
[orc-rt] Pass the reporting Session to Session error reporters. (#226622)
Session error reporters previously received only the Error, and had to
be supplied to the Session constructor, so they had no direct way to
refer to the Session reporting the error (e.g. to include its address in
a log message).
This commit changes ErrorReporterFn to take the reporting Session along
with the Error, and adds a Session::logErrors reporter that logs errors
via ORC_RT_LOG at Error level in the Session category. logErrors is only
available when Error-level logging is compiled in
(ORC_RT_LOG_ENABLED(Error)), so that it can't silently discard errors in
configurations where logging is compiled out. The noErrors and
AccumulateErrors test helpers are updated to work as both plain and
Session error reporters.
Replace cl::bits with cl::list. NFC (#226399)
cl::bits packs enum values into an unsigned, which would block migration
to TableGen based representation. The three users only test membership.
LLM-aided
[orc-rt] Move ORC_RT_LOG_ENABLED into config.h. (#226617)
ORC_RT_LOG_ENABLED depends only on the logging configuration values
defined in config.h, but was defined in Logging.h, so checking whether a
log level is compiled in required pulling in the logging backend's
headers (e.g. <os/log.h>).
This commit moves ORC_RT_LOG_ENABLED (and its level-token aliases) into
config.h, so that headers can check the logging configuration without
including Logging.h.
[CodeGen] Fix DetectDeadLanes for same-class copies with mismatched widths (#226556)
isCrossCopy() returned early whenever source and destination shared a
register class, treating the transfer as lane-for-lane. That is wrong
when a REG_SEQUENCE names a subregister slot narrower than the source
operand, e.g. a 32-bit vreg used as the lo16 piece of another 32-bit
vreg. Only take the early exit when the two widths match; otherwise fall
through to findCommonRegClass().
Fixes: ROCM-31212
