[pull] main from llvm:main - #1713
Merged
Merged
Conversation
…-uninitialized] (#213578) When compiled with `gcc`, the uninitialized `buf` triggers a warning. Since some external projects (e.g., `TPP-MLIR`) treat warnings as errors, the build fails. This patch initializes `buf` to eliminate the warning and prevent those build failures
Summary: This makes it easier to test a single function without opting-in to the whole `crt1.o` infra for `libc` that was originally intended to test existing CPU tests. Good for possible future unit tests, cheap tool that can launch a kernel. Only provides `foo(void)` kernels for now, can be improved.
…ranslation unit (#213720) CUFAddConstructor registered every CUF module variable mirrored in the GPU module, including variables brought in with USE from another translation unit. Without relocatable device code each translation unit has its own device module, so registering a declaration binds the host address to a module that does not contain the symbol, and a later cudaGetSymbolAddress fails with cudaErrorInvalidSymbol or cudaErrorSymbolNotFound. Under unified memory the same declaration is registered as host memory, overriding the defining unit's device registration. Skip globals that have no initializer and no body. Non-allocatable managed globals still register, since they go through a companion pointer local to the translation unit.
This is a follow-up to #153683 to support OpenMP compliant pointer-attachment in `declare_mappers` via `ATTACH`-style maps. In addition to enabling attach-style maps, we also need to propagate information about which map entries are for "pointee" data, i.e. have an "attach-ptr", and thus occupy a different storage block than the base variable for which the mapper is being generated. e.g. ```c S sa[10]; #pragma omp declare_mapper (default: S s) map (s.x, s.p[0:10]) #pragma omp target_enter_data map(sa) ``` The entry emitted for `s.p[0:10]` is for the pointee, i.e. it does not share storage with `s`. Mapper codegen needs to know that the entry for `sa[1].p[0:10]`, for example, is not a `MEMBER_OF` the map of `sa`, as it occupies its own storage and has its own ref-count tracking etc. Flang currently passes an unconditional `PreserveMemberofFlags` bool to the OMPIRBuilder function, which should eventually be propagated to using the per-entry information so that `map(s%x)` should get MEMBER_OF during mapper codgen, but `map(s%p(1:10))` should not. Currently, the per-entry MapInfo field is set to `false` for flang, so the change is a no-op for it. I don't have enough flang expertise/testing resources, so I'll let Andrew update Flang in follow-up changes. --------- Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>
cat and diff are the only two builtins that still spawn a subprocess: every cat/diff on a RUN line spawns a fresh Python interpreter, which dominates wall time given how small lit's typical inputs are. Run them in-process instead. The spawned-script path stays as a fallback for now. Removing it entirely is a follow-up PR. Output is byte-identical to the spawn path either way.
Reverts #209845 verified by local reverting unblock bots: https://lab.llvm.org/buildbot/#/builders/234/builds/1391 https://lab.llvm.org/buildbot/#/builders/10/builds/33193
**LROUND** with an i32 result on rv64 chose the libcall as `f64 ? LROUND_F64 : LROUND_F32`, so **fp128** fell through to the f32 libcall (lroundf) and lost precision. Use `RTLIB::getLROUND`, which maps **fp128** to lroundl.
SGPR lane spilling currently treats the WWM VGPR pool as all-or-nothing. This can fail compilation when the requested pool cannot be formed, even though scratch spilling or a smaller spillable pool could make progress. This PR lets ordinary SGPR spills fall back to scratch when the pool is empty and lets WWM register allocation use a nonempty partial pool. It keeps the full-pool requirement for strict WWM/WQM and explicit spill-carrier preallocation. The no-pool fallback is recorded in SIMachineFunctionInfo so frame lowering can provide enough emergency scavenging slots. The state is also serialized to preserve the behavior across MIR round trips. Fixes LCOMPILER-2542.
And replace the previous usages of llvm::PointerUnion with this new struct. This is currently just a refactoring but we'll later need to have a proper type for the previos typedef so we can pass it to emit* functions.
This PR adds the envar LIBOMPTARGET_JIT_SAVE_IMAGE_FILENAME to indicate that the device image generated after JIT compilation should be saved to a file.
) Both 1D vector.transfer_read / vector.transfer_write now lower to the scattered xegpu.load / xegpu.store path instead of xegpu.load_nd / xegpu.store_nd. We reserve the nd block path for rank ≥ 2 vectors. Cleanup Test: Where both RUN configurations (with and without --xevm-attach-target) produce identical IR for a function, the duplicated LOAD-ND/LOAD-GATHER (and STORE-ND/STORE-SCATTER) check blocks are collapsed into a single shared CHECK block. --------- Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
…ions (#204958) SimplifyCFG's foldCondBranchOnValueKnownInPredecessor redirects a predecessor edge around a block whose branch condition is known on that edge. When the bypassed block and the threading destination are on a common cycle, this can change the cycle structure and with it the dynamic instance of an uncontrolled convergent operation inside that cycle. On targets with branch divergence, this PR skips the candidate when the destination can reach an uncontrolled convergent call on a path back to the bypassed block. Operations using convergence control tokens are left alone. Fixes ROCM-26496.
When accessing a thread local variable, Clang generates a thread local wrapper function that checks if the variable has been initialized, and if it isn't, initializes it. This is a function call, so if this is within a funclet (i.e., within a `catchpad` or `cleanuppad`), it needs the funclet bundle argument, which was missing before. If it lacks a funclet argument, it will be considered invalid and removed in WinEHPrepare. Fixes emscripten-core/emscripten#27448.
…ring (#211947) Generalizes ContractionLowering in mlir/lib/Conversion/VectorToXeGPU/VectorToXeGPU.cpp so that (batched) N-D vector.contract ops lower to xegpu.dpas, not just plain 2D matmuls. --------- Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>
…212807) [RISCV] Reduce spill/reload pairs when Xqcilo extension is enabled Currently, `SelectAddrRegImm26` calls `SelectAddrFrameIndex` first, causing bare frame-index loads (offset 0) to select 48-bit loads/stores at ISel. Due to `AddedComplexity=2` on the QC48LdPat patterns, the wide opcode won over the standard LW/SW even though the resolved frame offset typically fits simm12. This led to more spills and reloads in functions which are under high register pressure because 48-bit loads and stores are not marked easily rematerializable. Also, simply adding 48-bit loads and stores to `isLoadFromStackSlot/isStoreToStackSlot` doesn't solve the regression for the multi call case and only by making Isel produce the plain 32/64-bit loads and store opcodes as the baseline does RA behave identically. Therefor this PR fixes the issue by: -Remove the `SelectAddrFrameIndex` call from SelectAddrRegImm26. Bare frame indices now fall through to standard LW/SW selection at ISel, where RA recognizes them as rematerializable stack loads. -Add post-RA promotion in `eliminateFrameIndex`: when a plain LW/SW has a resolved frame offset that exceeds simm12, promote the opcode to the corresponding 48-bit load/store opcode and fold the 26-bit offset directly. This preserves the intended large-offset optimization without affecting RA decisions. This solves the code size regression in high register pressure function introduced by PR #209315 Assisted by Claude
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to subscribe to this conversation on GitHub.
Already have an account?
Sign in.
15 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
See Commits and Changes for more details.
Created by
pull[bot] (v2.0.0-alpha.4)
Can you help keep this open source service alive? 💖 Please sponsor : )