[AMDGPU][Clang] add __builtin_amdgcn_av_(load|store)_b128 (#199176)
These builtins allow the program to request store-available and
load-visible accesses as described in #191246. Each of them takes a
__MEMORY_SCOPE_* operand that is then translated to target-specific
cache policy bits.
This patch was extracted from #172090.
Co-authored-by: macurtis-amd <macurtis at amd.com>
Assisted-by: Claude Opus 4.6
---------
Co-authored-by: macurtis-amd <macurtis at amd.com>
[mlir][linalg] Add sin, cos, tan to elementwise operations (#200950)
Add sin, cos, and tan as UnaryFn entries in the linalg dialect, enabling
their use via linalg.elementwise, named ops (linalg.sin, linalg.cos,
linalg.tan), and specialization from linalg.generic.
---------
Co-authored-by: Vinit Deodhar <vinitdeodhar at users.noreply.github.com>
Co-authored-by: Claude Opus 4.6 <noreply at anthropic.com>
Co-authored-by: Vinit Deodhar <vdeodhar at ah-vdeodhar-l.dhcp.mathworks.com>
[AMDGPU] SIFoldOperands: update BUNDLE header implicit use when folding
When folding an operand inside a BUNDLE, also rewrite the matching
implicit use on the bundle header. LiveVariables iterates a
MachineBasicBlock with the bundle-aware iterator and only inspects the
header, so without this update its kill flags go stale and a later
MachineVerifier run reports "Using a killed virtual register".
Co-Authored-By: Claude Opus 4 <noreply at anthropic.com>
[AMDGPU][GISel] Add register bank legalization rules for amdgcn_raw_buffer_load_async_lds (#201406)
Also amdgcn_struct_buffer_load_async_lds,
amdgcn_raw_ptr_buffer_load_async_lds, and
amdgcn_struct_ptr_buffer_load_async_lds.
[clang][bytecode] Fix an assertion failure in visitDtorCall() (#202507)
In `emitDestructionPop()`, we assert that the Descriptor has a
non-trivial dtor. Check this first here so we don't do all this work for
nothing.
[SLP] Update test against const-folding (#202532)
223ef1f3 ([IRBuilder] ConstFold unary intrinsics, #200496) made a lot of
test updates to SLPVectorizer. The tests were written a long time ago,
and it is unclear what their intent was, but at least update the one
test to replace constants with arguments, where the intent is clear.
[AMDGPU] Clamp load_monitor scope to minimum SCOPE_SE (#198245)
The load_monitor instructions monitor L2 cache lines and therefore
require at least SCOPE_SE to ensure the L2 cache is hit. The current
memory model requires the user to ensure that the specified scope is
such that it results in at least SCOPE_SE, otherwise the behaviour is
undefined. Instead, we now clamp the emitted scope at a minimum of
SCOPE_SE, so that the undefined behaviour is converted into a
performance loss instead.
Assisted-By: Claude Opus 4.6
[StringMap] Replace tombstone deletion with TAOCP 6.4 Algorithm R (#202103)
StringMap uses quadratic probing with lazy deletion: an erased entry
becomes a tombstone, a third bucket state alongside empty and live that
every find/insert must inspect.
Switch to linear probing with Knuth TAOCP 6.4 Algorithm R deletion,
similar to DenseMap #200595.
erase now relocates the following entries to close the hole. StringMap
buckets are pointers to heap-allocated entries, so only the pointers
(and the parallel hash array) move. References and pointers to entries
remain valid, but iterators are invalidated.
Depends on #202237 and #202520
Aided by Claude Opus 4.8
[InstCombine] Fix infinite combine loop in evaluateInDifferentType (#202572)
The implementation assumes that all original uses inside visited
instructions would get removed as part of changing the type. However,
this is not true for uses in select conditions, as only the value
operands change type in that case. Bail out if we encounter uses in
select conditions to avoid this.
Fixes https://github.com/llvm/llvm-project/issues/173148.
[AMDGPU][Clang] add __builtin_amdgcn_av_(load|store)_b128
These builtins allow the program to request store-available and load-visible
accesses as described in #191246. Each of them takes a __MEMORY_SCOPE_* operand
that is then translated to target-specific cache policy bits.
This patch was extracted from #172090.
Co-authored-by: macurtis-amd <macurtis at amd.com>
Assisted-by: Claude Opus 4.6z
[NewPM][AArch64][GlobalISel] Port AArch64PostLegalizerCombiner to NewPM (#194156)
Adds a standard porting.
Updates some (but not all) tests to verify the NewPM path is working.
[AMDGPU] Improve the description of asyncmark semantics
- The semantics of asyncmarks is now definded purely in terms of sequences,
without referring to the implementation.
- The examples incorrectly used (post)dominance. Fixed that with wording in
terms of asyncmark sequences.
[X86] Hygon Processors Initial enablement (#187622)
This patch adds initial support for several Hygon architectures.
The Hygon architectures include:
- c86-4g-m4
- c86-4g-m6
- c86-4g-m7
This patch includes:
- Added Hygon architectures CPU targets recognition in Clang and LLVM
- Added Hygon architectures to target parser and host CPU detection
- Updated compiler-rt CPU model detection for Hygon architectures
- Added Hygon architectures to various optimizer tests
- Added scheduler models for Hygon architectures CPU targets
Revert "[HLSL] Set visibility of cbuffer global variables to internal" (#202538)
Reverts llvm/llvm-project#200312
Breaks several buildbots, e.g.,
https://lab.llvm.org/buildbot/#/builders/203/builds/48531
Co-authored-by: Nikolas Klauser <nikolasklauser at berlin.de>
[mlir][bufferization] Drop TensorLikeType::getBufferType() (#201350)
Replace TensorLikeType::getBufferType() with
options.unknownTypeConverterFn() hook. Make the hook work with
tensor-like and buffer-like types (instead of builtins) to maintain the
same behaviour at the API boundary level and still allow user types to
be properly supported.
Historically, an attempt to support user types within the one-shot
bufferization framework was made. As part of it,
TensorLikeType::getBufferType() was introduced to allow user-provided
types to customize bufferization. However, the whole affair proved to be
overly complex: there is an interface with customization points for
user-provided tensors, and options-based (not sufficient) implementation
for builtin tensors. On top of this, there was always a
function-specific hook to customize function-level behaviour further. As
a result of this, users would need to implement two different mechanisms
on their end: interface implementation + option hooks.
[19 lines not shown]
[clang][bytecode] Fix shifting by negative IntAP values (#202505)
The negation of a negative value didn't necessarily result in a positive
value. Fix that by giving it one more bit of precision.
[OFFLOAD][L0] Add wait events for AsyncQueue memFill (#202287)
Fix an issue where memFill operations were not chained properly with respect prior operations.
[Loads] Migrate isDereferenceable APIs to SimplifyQuery (#202553)
These take the usual set of analysis parameters, so we can encapsulate
them using SimplifyQuery.
[Clang][HIP] Include `__clang_cuda_math_forward_declares.h` before `<cmath>` (#201563)
In HIP, `constexpr` functions are treated as both `__host__` and
`__device__`.
A new version of the MS STL shipped with the build tools version
14.51.36231 has `constexpr` definitions for some `cmath` functions when
the
compiler in use is Clang (this gets worse when C++23 is in use).
These definitions conflict with the `__device__` declarations we provide
in the header wrappers.
There is a workaround for this: We do not mark `constexpr`
functions [_that are defined in a system
header_](https://github.com/llvm/llvm-project/blob/03127a03860b9d8cb440fe8f51c00647f45eb8be/clang/lib/Sema/SemaCUDA.cpp#L877)
as
`__host__` and `__device__` if there is a previous `__device__`
declaration.
[14 lines not shown]
