[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
[AMDGPU] Measure MFMA overwrite hazards at each instruction
Apply previously established processing to:
- VALU overwriting an MFMA result
- VALU overwriting a register an MFMA took as srcC
AI-assisted.
[AMDGPU] Measure MFMA read hazards at each producer
Introduce more sophisticated traversal to avoid the following traps:
- order-dependent traversal and discarding seen BBs despite shorter path
- mis-matching distance and window of different producers
Record the best distance per BB instead of a visited flag and sweep the
arrivals in nondecreasing distance (bucket queue). This pairs producers
with their actual distance to a consumer in one go.
Fixed scenarios:
- MFMA reading an MFMA result as srcA, srcB or srcC
- VALU, memory or export instruction reading an MFMA result
rewrite-vgpr-mfma-to-agpr.ll gains an s_nop 2: a 4-pass XDL write that
partially overlaps the srcC read two slots later requires five wait
states, and none were emitted because the nearest producer wrote the
register in full.
AI-assisted.
[AMDGPU][NFC] Extract the MFMA read-window calculation
Move the wait states a consumer needs before reading an MFMA result out
of checkMAIHazards90A into getMFMAReadWaitStates, taking the producer as
an argument, so a caller can ask about a specific producer. The partial
srcC overlap half moves into getMFMAOverlappedSrcCWaitStates. The caller
passes the producer the walk recorded, so nothing changes.
AI-assisted.
