InstCombine: Stop applying nofpclass from use nofpclass attribute (#183835)
Functionally reverts a80d4329ce96856a02bd279c800c3d08619da4c9, with new
test.
This should be applied somewhere, but this is the wrong place.
Fixes regression reported after #182444
[lldb/test] Skip TestDelayInitDependency on remote platforms (#183885)
This test exercises macOS-specific linker functionality (-delay_library)
and uses a hardcoded local working directory for the launch info. It
should not run against a remote platform where neither condition holds.
Signed-off-by: Med Ismail Bennani <ismail at bennani.ma>
[clang] Backport: fix transformation of substituted constant template parameters of partial specializations
This fixes a helper so it implements retrieval of the argument replaced
for a template parameter for partial spcializations.
This was left out of the original patch, since it's quite hard to actually test.
This helper implements the retrieval for variable templates, but only for
completeness sake, as no current users rely on this, and I don't think a similar
test case is possible to implement with variable templates.
This fixes a regression introduced in #161029 which will be backported to llvm-22,
so there are no release notes.
Backport from #183348
Fixes #181062
Fixes #181410
[AMDGPU] Add structural stall heuristic to scheduling strategies
Implements a structural stall heuristic that considers both resource
hazards and latency constraints when selecting instructions. In coexec,
this changes the pending queue from a binary “not ready to issue”
distinction into part of a unified candidate comparison. Pending
instructions still identify structural stalls in the current cycle, but
they are now evaluated directly against available instructions by stall
cost, making the heuristics both more intuitive and more expressive.
- Add getStructuralStallCycles() to GCNSchedStrategy that computes the
number of cycles an instruction must wait due to:
- Resource conflicts on unbuffered resources (from the SchedModel)
- Sequence-dependent hazards (from GCNHazardRecognizer)
- Add getHazardWaitStates() to GCNHazardRecognizer that returns the number
of wait states until all hazards for an instruction are resolved,
providing cycle-accurate hazard information for scheduling heuristics.
[AMDGPU] Add ML-oriented coexec scheduler selection and queue handling
This patch adds the initial coexec scheduler scaffold for machine
learning workloads on gfx1250.
It introduces function and module-level controls for selecting the
AMDGPU preRA and postRA schedulers, including an `amdgpu-workload-type`
module flag that maps ML workloads to coexec preRA scheduling and a nop
postRA scheduler by default.
It also updates the coexec scheduler to use a simplified top-down
candidate selection path that considers both available and pending
queues through a single flow, setting up follow-on heuristic work.
[Driver] Add -Wa,--reloc-section-sym= to control section symbol conversion (#183472)
Wire the llvm-mc --reloc-section-sym={all,internal,none} option through
the clang driver (-Wa,--reloc-section-sym=) and cc1as
(--reloc-section-sym=). The option is only valid for ELF targets.
GNU Assembler will add the option as well.
[AMDGPU] Fix piggybacking after commute in AMDGPULowerVGPREncoding (#183778)
After successfully commuting an instruction to be compatible with the
current VGPR MSB mode, update CurrentMode with the commuted
instruction's mode requirements. This locks in the mode bits the
commuted instruction relies on, preventing later instructions from
piggybacking and corrupting those bits.
Without this fix, a subsequent instruction needing a different mode
could piggyback onto the preceding s_set_vgpr_msb and change mode bits
that the commuted instruction depends on. For example, a nullopt src1
position (treated as 0) could be overwritten to a different value,
causing incorrect register encoding for the commuted instruction.
The fix still allows compatible piggybacking - instructions that only
add new mode bits without changing existing ones can still piggyback.
AArch64: Replace @plt/%gotpcrel in data directives with %pltpcrel %gotpcrel (#155776)
Similar to #132569 for RISC-V, replace the unofficial `@plt` and
`@gotpcrel` relocation specifiers, currently only used by clang
-fexperimental-relative-c++-abi-vtables, with %pltpcrel %gotpcrel. The
syntax is not used in humand-written assembly code, and is not supported
by GNU assembler.
Also replace the recent `@funcinit` with `%funcinit(x)`.
[CodeGen] Allow `-enable-ext-tsp-block-placement` and `-apply-ext-tsp-for-size` passed together (#183642)
Currently, the asserts fires when both `UseExtTspForPerf` and
`UseExtTspForSize` are true on a given function.
Ideally, we should allow `-enable-ext-tsp-block-placement` and
`-apply-ext-tsp-for-size` passed together, meaning run the block
placement for performance on hot functions, while run the placement for
size on cold functions.
The diff makes `UseExtTspForPerf` and `UseExtTspForSize` mutually
exclusive per-function: functions with the `OptForSize` attribute use
ext-tsp block placement for size, while the others use ext-tsp block
placement for perf.
Co-authored-by: Sharon Xu <sharonxu at fb.com>
[CIR] Use `-verify` on clang/test/CIR/CodeGenHLSL/matrix-element-expr-load.hlsl (#182817)
Update clang/test/CIR/CodeGenHLSL/matrix-element-expr-load.hlsl to use
`-verify` with expected CIR NYI diagnostics.
[CIR] Infrastructure and MemorySpaceAttrInterface for Address Spaces (#179073)
Related: https://github.com/llvm/llvm-project/issues/175871,https://github.com/issues/assigned?issue=llvm%7Cllvm-project%7C179278,https://github.com/issues/assigned?issue=llvm%7Cllvm-project%7C160386
- Introducing the LangAddressSpace enum with offload address space kinds
(offload_private, offload_local, offload_global, offload_constant,
offload_generic) and the LangAddressSpaceAttr attribute.
- Generalizes CIR AS attributes as MemorySpaceAttrInterface and Attaches
it to `PointerType`. Includes test coverage for valid IR roundtrips and
invalid address space parsing.
This starts a series of patches with the purpose of bringing complete
address spaces support features for CIR. Most of the test coverage is
provided in subsequent patches further down the stack. note that most of
these patches are based on: https://github.com/llvm/clangir/pull/1986
[VPlan] Don't adjust trip count for DataAndControlFlowWithoutRuntimeCheck (#183729)
Previously, the canonical IV increment may have overflowed to a non-zero
value due to vscale being a non power-of-two. So we used to emit a
runtime check for this.
If you didn't want the runtime check,
DataAndControlFlowWithoutRuntimeCheck skipped it and instead tweaked the
trip count so it wouldn't overflow.
However #144963 stopped the check from ever being emitted because vscale
is always a power-of-two on AArch64 and RISC-V, so it never overflowed
to a non-zero value. And in #183292 the code to emit the check was
removed. But we never restored the trip count back to normal when the
target's vscale was a power-of-two.
Now that vscale is always a power-of-two, this PR avoids adjusting it. A
follow up NFC can then remove DataAndControlFlowWithoutRuntimeCheck.
Clang: Deprecate float support from __builtin_elementwise_max (#180885)
Now we have
__builtin_elementwise_maxnum
__builtin_elementwise_maximum
__builtin_elementwise_maximumnum
[libc][math][c23] implement C23 `acospif` math function (#183661)
Implementing C23 `acospi` math function for single-precision with the
header-only approach that is followed since #147386
[libc][math] Refactor floor family to header-only (#182194)
Refactors the floor math family to be header-only.
Closes https://github.com/llvm/llvm-project/issues/182193
Target Functions:
- floor
- floorbf16
- floorf
- floorf128
- floorf16
- floorl
[AMDGPU] Remove extra pipes from load-saddr-offset-imm.ll (#183874)
This test uses opt to run instcombin and then pipes that into llc which
has its output piped into FileCheck. Before this patch, the test also
piped in the source file into llc as well, which caused issues with a
downstream test executor that executes the lines in bash. However, these
extra pipes don't make sense anyways, so remove them.
[clang] NFC: remove unused / untested workaround in pack deduction
This snippet was part of what was introduced in 130cc445e46836b28defdce03b1adfdb16ddcf41
However, none of the existing tests require it, including the tests added in
that commit.
One of those tests had a FIXME which was fixed when we switched
frelaxed-template-template-args on by default as well.
[CUDA] Allow `extern __shared__` on non-array types
NVCC allows `extern __shared__` on any type, not just incomplete arrays.
This is commonly used in CUDA libraries like NCCL to overlay a struct on
dynamically-allocated shared memory:
extern __shared__ ncclShmemData ncclShmem;
Previously, Clang rejected this with a hard error and did not add
`CUDASharedAttr` to the VarDecl. This caused a cascade: `IdentifyTarget()`
classified the variable as host-side, and any device code referencing it
got a spurious "reference to __host__ variable in __device__ function"
error.
Downgrade the error to a default-ignored warning (`-Wcuda-extern-shared`)
and always add `CUDASharedAttr` so the variable is correctly classified as
device-side. The old `err_cuda_extern_shared` is preserved for potential
future use.
[AMDGPU][SIInsertWaitcnts] Move VCCZ workaround code out of the way (#182619)
This is a cleanup patch that moves the VCCZ specific workaround code
from `SIInsertWaitcnts::insertWaitcntInBlock()` to a separate class and
refactors it a bit to make it easier to read.
The end result is a simpler `insertWaitcntInBlock()`.
Should be NFC.
[CIR][NFC] Move some builtin tests to the CodeGenBuitins folder (#183607)
This moves a few tests that were created in the wrong location. Also
changes the names of some test files to maintain consistency.
Fix profile metadata propagation in InstCombine select folding
Propagate profile metadata when folding select instructions with logical AND/OR conditions and when canonicalizing SPF to intrinsics. This fixes profile verification failures in Transforms/InstCombine/select-and-or.ll.
[SLP]Fix operand reordering when estimating profitability of operands
Need to swap operand for a single instruction, not for the the same lane
of the first and second instruction in the list
