These concepts don't exist for MCRegister.
I think there is a need for virtual registers in MCRegister for NVPTX,
SPIR-V and WebAssembly, but those should not be confused with Register's
virtual register. I will try to make a separate proposal for that with a
real interface.
They were instantiated with only a single type and union-members
themselves. By putting the types directly into a union, they are still
left uninitialized by default.
Replace binary of of two reductions with one reduction of the binary op
applied to vectors. For example:
```
%v0_red = tail call i32 @llvm.vector.reduce.add.v16i32(<16 x i32> %v0)
%v1_red = tail call i32 @llvm.vector.reduce.add.v16i32(<16 x i32> %v1)
%res = add i32 %v0_red, %v1_red
```
gets transformed to:
```
%1 = add <16 x i32> %v0, %v1
%res = call i32 @llvm.vector.reduce.add.v16i32(<16 x i32> %1)
```
Common code has been unified and generalized.
Original commit: 123dca9b56e1359d8ec7771ea3bd0afd4b1ea6af
Previously reverted due to accidentally merged incompletely. The issue has
been addressed by restoring missing code.
This reverts commit 123dca9b56e1359d8ec7771ea3bd0afd4b1ea6af.
This breaks building on macOS with clang and multiple build bots,
including https://lab.llvm.org/buildbot/#/builders/175/builds/13585
llvm-project/llvm/lib/Transforms/Utils/SimplifyCFG.cpp: In function ‘bool sinkCommonCodeFromPredecessors(llvm::BasicBlock*, llvm::DomTreeUpdater*)’:
/b/ml-opt-devrel-x86-64-b1/llvm-project/llvm/lib/Transforms/Utils/SimplifyCFG.cpp:2503:3: error: reference to ‘LockstepReverseIterator’ is ambiguous
2503 | LockstepReverseIterator<true> LRI(UnconditionalPreds);
| ^~~~~~~~~~~~~~~~~~~~~~~
I noticed that DIEInteger::BestForm used a cast to char:
if ((char)Int == SignedInt)
If 'char' happens to be unsigned, this will not behave correctly. Then I
also noticed that this code assumes the size of 'short' and 'int'.
This patch changes this code to use more precise types. No functional
change should be visible on ordinary hosts.
Common code has been unified and generalized. Not sure if it may be
worth to generalize this further, since it looks closely tied to Blocks
(might make sense to rename it in `LockstepReverseInstructionIterator`).
`JSONScopedPrinter` has a `std::unique_ptr<DelimitedScope>` member and
defaulted constructor argument, so it needs a complete type. This
resolves one of the many build errors with C++23 using Clang.
DICompositeType has an attribute representing the name of a type, but
currently it isn't possible to set this for array types via the
DIBuilder method. This patch adds a new overload of
DIBuilder::createArrayType that allows "full" construction of an array
type. This is useful for Ada, where arrays are a bit more first-class
than C.
Prior to gcc version 9, the `__PRETTY_FUNCTION__` macro was not declared
constexpr. In that case, don't declare this as constexpr, and switch the
static asserts to runtime asserts.
Verified this should work on all supported compilers:
https://godbolt.org/z/T77rvPW5z
Followup to #127893 / 8a39214b7e6e79e65135b7dbcc98b427a032d5fb
This is a followup to 003a721c1c9e3a99d6d0c1a6755443b260235537, which we
noticed increased binary size a small but noticable amount. The increase
seems to be due to code size from each `llvm::getTypeName<T>` wrapper,
which adds up if there are a lot of types.
The original motivation was to improve runtime and avoid `getTypeName`
being recomputed each time. The implementation is simple enough that we
can make it fully constexpr, which addresses the size increase, but also
further improves runtime: we directly reference the data instead of
jumping through a guard variable to see if we've computed it already.
Callee saved registers should always be phyiscal registers. They are
often passed directly to other functions that take MCRegister like
getMinimalPhysRegClass or TargetRegisterClass::contains.
Unfortunately, sometimes the MCRegister is compared to a Register which
gave an ambiguous comparison error when the MCRegister is on the LHS.
Adding a MCRegister==Register comparison operator created more ambiguous
comparison errors elsewhere. These cases were usually comparing against
a base or frame pointer register that is a physical register in a
Register. For those I added an explicit conversion of Register to
MCRegister to fix the error.
This patch implements the check for not allowing re-scheduling of
instructions that have already been scheduled in a scheduling bundle.
Rescheduling should only happen if the instructions were temporarily
scheduled in singleton bundles during a previous call to
`trySchedule()`.
Up until now the generation of vector instructions was taking place
during the top-down post-order traversal of vectorizeRec(). The issue
with this approach is that the vector instructions emitted during the
traversal can be reordered by the scheduler, making it challenging to
place them without breaking the def-before-uses rule.
With this patch we separate the vectorization decisions (done in
`vectorizeRec()`) from the code generation phase (`emitVectors()`). The
vectorization decisions are stored in the `Actions` vector and are used
by `emitVectors()` to drive code generation.
Use nonstatic member instead. This requires explicit conversions, but
many will go away as we continue converting unsigned to Register.
In a few places where it was simple, I changed unsigned to Register.
Change getDependenceDistanceStrideAndSize to scale strides by
TypeByteSize, scaling the returned CommonStride and MaxStride. Even
though there is a seemingly-functional change of setting CommonStride
when scaled strides are equal, it ends up being a non-functional change
due to aggressive HasSameSize checking.
Shaders that use the llvm.amdgcn.init.whole.wave intrinsic need to
explicitly preserve the inactive lanes of VGPRs of interest by adding
them as dummy arguments. The code usually looks something like this:
```
define amdgcn_cs_chain void f(active vgpr args..., i32 %inactive.vgpr1, ..., i32 %inactive.vgprN) {
entry:
%c = call i1 @llvm.amdgcn.init.whole.wave()
br i1 %c, label %shader, label %tail
shader:
[...]
tail:
%inactive.vgpr.arg1 = phi i32 [ %inactive.vgpr1, %entry], [poison, %shader]
[...]
; %inactive.vgpr* then get passed into a llvm.amdgcn.cs.chain call
```
Unfortunately, this kind of phi node will get optimized away and the
backend won't be able to figure out that it's ok to use the active lanes
of `%inactive.vgpr*` inside `shader`.
This patch fixes the issue by introducing a llvm.amdgcn.dead intrinsic,
whose result can be used as a PHI operand instead of the poison. This
will be selected to an IMPLICIT_DEF, which the backend can work with.
At the moment, the llvm.amdgcn.dead intrinsic works only on i32 values.
Support for other types can be added later if needed.
Similar to #117309.
The advisor and logger are accessed through the provider, which is
served by the new PM. Legacy PM forwards calls to the provider.
New PM is a machine function analysis that lazily initializes the
provider.
This reverts commit 36eaf0daf5d6dd665d7c7a9ec38ea22f27709fed.
This is not a sound approach to dealing with this instruction change.
The new behavior is a different opcode pair, not a modifier on the
existing opcode.
A non-escaping global should never alias with non-pointer values
and `ptr null`. This should improve disambiguation of global
pointers with relation to Flang runtime calls (given that
`nosync nocallback` attributes are properly set).
It also seems to be an obvious improvement with little overhead.
This patch implements a new subclass of the Value class used for Sandbox
IR Values that we don't support, like metadata or inline asm. The goal
is to never have null sandboxir::Value objects, because this is not the
expected behavior.
For targets that support IEEE fminimum_num/fmaximum_num, the
corresponding *_min_num_fXY/*_max_num_fXY instructions themselves
already did the canonicalization for the inputs. As a result, we do not
need to explicitly canonicalize the inputs for fminnum/fmaxnum.
This patch adds intrinsics for tcgen05.cp and
tcgen05.shift instructions.
lit tests are added and verified with a
ptxas-12.8 executable.
Docs are updated in the NVPTXUsage.rst file.
Signed-off-by: Durgadoss R <durgadossr@nvidia.com>
gfx940 and gfx941 are no longer supported. This is one of a series of
PRs to remove them from the code base.
This PR removes all non-documentation occurrences of gfx940/gfx941 from
the llvm directory, and the remaining occurrences in clang.
Documentation changes will follow.
For SWDEV-512631
The standard libcalls for half to float and float to half conversion are
__extendhfsf2 and __truncsfhf2. However, LLVM currently uses
__gnu_h2f_ieee and __gnu_f2h_ieee instead. As far as I can tell, these
libcalls are an ARM-ism and only provided by libgcc on that platform.
compiler-rt always provides both libcalls.
Use the standard libcalls by default, and only use the __gnu libcalls on
ARM.
Adds `findByUse` which takes a `llvm::Value` from a use and resolves it
(as best as possible) back to the creation of that resource.
It may return multiple ResourceBindingInfo if the use comes from
branched control flow.
Fixes#125746
In a particular scenario (see test) we used to insert scheduled
instructions into the ready list. This patch fixes this by fixing the
trimSchedule() function.
This patch implements a small region pass that saves the context's
state. The patch is now used in the default pipeline to save the context
state instead of the hard-coded call to Context::save().
The concept behind this is that the passes themselves should not have to
do the actual saving/restoring of the IR state, because that would make
it challenging to reorder them in the pipeline. Having separate
save/restore passes makes the transformation passes more composable as
parts of arbitrary pipelines.
Attempting to pass a `ptr addrspace(7)` to functions that take `ptr`
arguments produces undesirable `addrspacecast(addrspacecast(p8 x to p7)
to p0) => addrspacecast(p8 x to p0)` folds. This results in illegal GEP
operations on buffer resources, which can't be GEP'd. (However, note
that, while unimplemneted, addressspacecast from ptr addrspace(7) to ptr
is legal - it's just an effective address computation)
To resolve this problem, and thus prevent illegal
`getelementptr T, ptr addrspace(8) %x, ...` s from being produces, this
commit extends amdgcn.make.buffer.rsrc to also be variadic in its result
type, auto-upgrading old manglings.
The logic for handling a make.buffer.rsrc in instruction selection
remains untouched and expects the output type to be a ptr addrspace(8),
as does the Clang lowering for its builtin (the pointer-to-pointer
version might want a different name in clang). LowerBufferFatPointers
has been updated to lower
amdgcn.make.buffer.rsrc.p7.p* to amdgcn.make.buffer.rsrc.p8.p* .
This'll also make exposing buffer fat pointers in Clang easier, since
you don't have to cast between a `__amdgcn_rsrc_t` and a pointer.
