Previously the OpenCL address space overloads of remquo would call into
the one and only 'private' CLC remquo. This was an outlier compared with
the other pointer-argumented maths builtins.
This commit moves the definitions of all address space overloads to the
CLC library to give more control over each address space to CLC
implementers.
There are some minor changes to the generated bytecode but it's simply
moving IR instructions around.
The conditional variants of SMSTART/STOP currently take the current
PStateSM as a variadic value. This is not supported by the verification
added in #140472 (which requires variadic values to be of type Register
or RegisterMask), so this patch splits the the conditional variants into
new `COND_` nodes, where these extra parameters are fixed arguments.
Suggested in
https://github.com/llvm/llvm-project/pull/140472#discussion_r2094635066
Part of #140472.
After the introduction of OpAsmAttr/TypeInterface in #121187#124721,
the documentation for them could be updated along side the doc for
OpAsmDialectInterface.
#127993 changed the trailing digit behavior for alias name generation.
As best as I can see, all NVPTX architectures support the generic
address space.
I note there's a FIXME in the target's address space map about 'generic'
still having to be added to the target but we haven't observed any
issues with it downstream. The generic address space is mapped to the
same target address space as default/private (0), but this isn't
necessarily a problem for users.
The Chromium bug tracker is in an archived state. The Security Response
Group has preemptively created llvm-project GitHub issues with PDF
copies of the Chromium issues should the repository become inaccessible.
* Add URLs for redirects from
https://bugs.chromium.org/p/llvm/issues/detail?id=X to
https://issuetracker.google.com/issues/y
* Add URLs to llvm-project archive issues.
* Add an explanation of archive use.
- Deleted unused includes
- Deleted useless braces
- Modernized tests to use `CHECK-MESSAGES-NOT` and `CHECK-FIXES-NOT` for
better readability and maintainability
Currently VectorCombine can scalarize vector compares and binary ops.
This extends it to also scalarize binary-op like intrinsics like umax,
minnum etc.
The motivation behind this is to scalarize more intrinsics in
VectorCombine rather than in DAGCombine, so we can sink splats across
basic blocks: see #137786
This currently has very little effect on generated code because
InstCombine doesn't yet canonicalize binary intrinsics where one operand
is a constant into the form that VectorCombine expects, i.e. `binop
(shuffle insert) const --> shuffle (binop insert const)`. The plan is to
land this first and then in a subsequent patch teach InstCombine to do
the canonicalization to avoid regressions in the meantime.
This uses `isTriviallyVectorizable` to determine whether or not an
intrinsic is safe to scalarize. There's also `isTriviallyScalarizable`,
but this seems more geared towards the Scalarizer pass and includes
intrinsics with multiple return values.
It also only handles intrinsics with two operands with the same type as
the return type. In the future we would generalize this to handle
arbitrary numbers of operands, including unary operators too, e.g. fneg
or fma, as well as different operand types, e.g. powi or scmp
try_emplace can default-construct values, so we do not need to do so
on our own. Plus, try_emplace(Key) is much shorter than
insert(std::make_pair(Key, Value()).
When rewriting multiple CompositeInserts to CompositeConstruct, we need
to know the number of elements of the result type. However, we cannot
query the number of elements for cooperative matrix types.
Skipping only over the first results leads to the curCodeIt pointing to
the wrong location in the bytecode, causing the execution to continue
with a wrong instruction after the Constraint/Rewrite.
Signed-off-by: Rickert, Jonas <Jonas.Rickert@amd.com>
[analyzer][NFC] Move PrettyStackTraceLocationContext into
dispatchWorkItem
This change helps with ensuring that the abstract machine call stack is
only dumped exactly once no matter what checker callback we have the
crash in.
Note that `check::EndAnalysis` callbacks are resolved outside of
`dispatchWorkItem`, but that's the only checker callback that is outside
of `dispatchWorkItem`.
CPP-6476
Try to do as few checks as possible. Check for builtin types only once,
then look at the BuiltinType Kind. For integers, we cache the int and
long size, since those are used a lot and the ASTContext::getIntWidth()
call is costly.
Fix the handling of the `instructionOffset` parameter, which resulted in
always returning the wrong disassembly because VSCode always uses
`instructionOffset = -50` and expects 50 instructions before the given
address, instead of 50 bytes before
This pull request fixes a couple of unhandled situations in DWARF input
leading to a crash. Specifically,
- If the DWARF input contains a declaration of a C variadic function
(where `...` translates to `DW_TAG_unspecified_parameters`), which is
then followed by a definition, `llvm_unreachable()` is hit in
`LVScope::addMissingElements()`. This is only visible in Debug builds.
- Parsing of instructions in `LVBinaryReader::createInstructions()` does
not check whether `Offset` lies within the `Bytes` ArrayRef. A specially
crafted DWARF input can lead to this condition.
Motivation:
Currently, the NVVMOps are not verified against the supported SM
architectures. This can manifest as an ISel failure in the NVPTX LLVM
backend during CodeGen to PTX ISA. This PR addresses this issue by
adding verifier checks for Target-SM architectures in the NVVM Dialect
itself, thereby catching the errors early on.
Summary:
* Parametric traits named `NVVMRequiresSM` and `NVVMRequiresSMa` are
added to facilitate the version checks for typical and arch-accelerated
versions respectively.
* These traits can be attached to any NVVM Op to enable the checks for
the particular Op. (example shown below)
* An attribute interface called named `TargetAttrVerifyInterface` is
added to the GPU dialect which any target attribute seeking to perform
target-verification on the module can implement.
* The checks are performed by the `NVVMTargetAttr` (implementing the
`TargetAttrVerifyInterface` interface) when called from the GPU module
verifier where it walks through the module and performs the checks for
Ops with the `NVVMRequiresSM` traits.
* A few Ops in `NVVMOps.td` have been updated to serve as examples.
Example Usage:
```
def NVVM_ReduxOp : NVVM_Op<"redux.sync"> {...}
----> def NVVM_ReduxOp : NVVM_Op<"redux.sync", [NVVMRequiresSM<80>]> {...}
def NVVM_WgmmaFenceAlignedOp : NVVM_Op<"wgmma.fence.aligned"> {...}
----> def NVVM_WgmmaFenceAlignedOp : NVVM_Op<"wgmma.fence.aligned", [NVVMRequiresSMa<[90]>]> {...}
```
---------
Co-authored-by: Guray Ozen <guray.ozen@gmail.com>
The generic GPU barrier implementation checked if it was the main thread
in generic mode to identify single threaded regions. This doesn't work
since inside of a non-active (=sequential) parallel, that thread becomes
the main thread of a team, and is not the main thread in generic mode.
At least that is the implementation of the APIs today.
To identify single threaded regions we now check the team size
explicitly.
This exposed three other issues; one is, for now, expected and not a
bug, the second one is a bug and has a FIXME in the
single_threaded_for_barrier_hang_1.c file, and the final one is also
benign as described in the end.
The non-bug issue comes up if we ever initialize a thread state.
Afterwards we will never run any region in parallel. This is a little
conservative, but I guess thread states are really bad for performance
anyway.
The bug comes up if we optimize single_threaded_for_barrier_hang_1 and
execute it in Generic-SPMD mode. For some reason we loose all the
updates to b. This looks very much like a compiler bug, but could also
be another logic issue in the runtime. Needs to be investigated.
Issue number 3 comes up if we have nested parallels inside of a target
region. The clang SPMD-check logic gets confused, determines SPMD (which
is fine) but picks an unreasonable thread count. This is all benign, I
think, just weird:
```
#pragma omp target teams
#pragma omp parallel num_threads(64)
#pragma omp parallel num_threads(10)
{}
```
Was launched with 10 threads, not 64.
When using InitChecker with VerifyOnly, we create a new designated
initializer to handle anonymous fields. However in the last call to
CheckDesignatedInitializer, the subinitializer isn't properly preserved
but it gets overwritten by the cloned one. Which causes the initializer
to reference the dependent field, breaking assumptions when we
initialize the instantiated specialization.
Fixes https://github.com/llvm/llvm-project/issues/67173
I got a bug report where a pedantic DAP client complains about getting
two "new" module events for the same UUID. This is caused by the dyld
transition from the on-disk dyld to the shared cache dyld, which share
the same UUID. The transition is not generating an unloaded event
(because we're not really unloading dyld) but we do get a loaded event
(because the load address changed). This PR fixes the issue by relying
on the modules set as the source of truth instead of relying on the
event type.
Currently the liveness analysis always marks operands yielded in regions
that aren't classified as `RegionBranchOpInterface` or
`CallableOpInterface` as non-live. Examples for these ops include
linalg.generic (with `linalg.yield` as terminator) or gpu ops (with
`gpu.yield` as terminator).
This in turn makes the `remove-dead-values` pass always incorrectly
remove the bodies of these ops, leading to invalid IR. Because these ops
define their own semantics, I have conservatively marked all operands of
these yield ops to be live.
When using -no-pie without a SECTIONS command, the linker uses the
target's default image base. If -Ttext= or --section-start specifies an
output section address below this base, the result is likely unintended.
LLD will give a diagnostic (#140187) and may change the behavior in the future.
It's good to set an explicit image base to avoid relying on its current
behavior. BOLT doesn't seem to care whether a PT_PHDR segment is
present.
Pull Request: https://github.com/llvm/llvm-project/pull/140570
I couldn't find an existing way to pass -mcpu=lime1 equivalent to LTO
codegen.
This commit would privide one. With this commit, you can do so by
passing
`-mllvm -mcpu=lime1` to wasm-ld.
