- Sep 22, 2023
-
-
Kiran Chandramohan authored
-
Anatoly Trosinenko authored
To simplify handling PAuth in the machine outliner, introduce a separate AArch64PointerAuth pass that is executed after both Prologue/Epilogue Inserter and Machine Outliner passes. After moving to AArch64PointerAuth, signLR and authenticateLR are not used outside of their class anymore, so make them private and simplify accordingly. The new pass is added via AArch64PassConfig::addPostBBSections(), so that it can change the code size before branch relaxation occurs. AArch64BranchTargets is placed there too, so it can take into account any PACI(A|B)SP instructions and not excessively add BTIs at the start of functions. Reviewed By: tmatheson Differential Revision: https://reviews.llvm.org/D159357
-
Luke Lau authored
-
Haojian Wu authored
-
David Green authored
This fills out some extra cases for sin/cos testing for various types under Global ISel, which seem to all do OK. The existing tests in sincospow-vector-expansion.ll can be removed, as they are now covered elsewhere.
-
Kiran Chandramohan authored
-
Mirko Brkusanin authored
-
Nikita Popov authored
-
Nikita Popov authored
This is a followup to #66988. The implementation there did not account for the possibility of the catch object frame index referrring to a fixed object, which is the case on win64.
-
Benjamin Maxwell authored
This attribute makes the `enable_arm_streaming` pass ignore a function (i.e. not add the enable streaming/za attributes). The main use case for this is to prevent helper functions within tests being made streaming functions.
-
qcolombet authored
Prior to this patch, `GeneralizeOuterUnitDimsUnPackOpPattern` would assert that we cannot create a `tensor.empty` operation with dynamic shapes. The problem stems from the fact that we were not using the right builder for the `tensor.empty` operation. Indeed, each dynamic dim needs to be specified by an input variable. Simply provide the dynamic dimensions to the `tensor.empty` builder to fix that.
-
Florian Hahn authored
After 4a5bcbd5, switch instructions can now be handled in a straight-forward manner by adding (ICMP_EQ, ConditionVal, CaseVal) for te successor blocks per case.
-
Ivan Kosarev authored
We don't seem to have a use for the -amdgpu-keep-16-bit-reg-suffixes option anymore. Was introduced in <https://reviews.llvm.org/D79435>. Reviewed By: Joe_Nash, foad Differential Revision: https://reviews.llvm.org/D156102
-
Florian Hahn authored
Shorten the types used to i8 for cheaper verification and add test case where 2 cases have the same destination, as suggested in #67061.
-
Simon Pilgrim authored
[DAG] getNode() - remove oneuse limit from (zext (trunc (assertzext x))) -> (assertzext x) fold (REAPPLIED) Noticed on D159533 and I've finally dealt with the x86 regressions - MatchingStackOffset wasn't peeking through AssertZext nodes while trying to find CopyFromReg/Load sources, it was only removing them if they were part of a (trunc (assertzext x)) pattern. Reapplied after being reverted at 4389252c - which should be addressed by D159537 / 6d267999
-
Alcaro authored
-
Ivan Kosarev authored
The existing fake True16 instructions using 32-bit VGPRs are supposed to co-exist with real ones until all the necessary True16 functionality is implemented and relevant tests are updated. Reviewed By: arsenm, Joe_Nash Differential Revision: https://reviews.llvm.org/D156101
-
Nikita Popov authored
The write to the SEH catch object happens before cleanuppads are executed, while the first reference to the object will typically be in a catchpad. If we make use of first-use analysis, we may end up allocating an alloca used inside the cleanuppad and the catch object at the same stack offset, which would be incorrect. https://reviews.llvm.org/D86673 was a previous attempt to fix it. It used the heuristic "a slot loaded in a WinEH pad and never written" to detect catch objects. However, because it checks for more than one load (while probably more than zero was intended), the fix does not actually work. The general approach also seems dubious to me, so this patch reverts that change entirely, and instead marks all catch object slots as conservative (i.e. excluded from first-use analysis) based on the WinEHFuncInfo. As far as I can tell we don't need any heuristics here, we know exactly which slots are affected. Fixes https://github.com/llvm/llvm-project/issues/66984.
-
Ivan Kosarev authored
Real True16 instructions are as they are defined in the ISA. Fake True16 instructions are identical to real ones except that they take 32-bit registers as operands and always use their low halves. Reviewed By: Joe_Nash Differential Revision: https://reviews.llvm.org/D156100
-
Guray Ozen authored
This work introduces a new operation called `warpgroup.mma` to the NVGPU dialect of MLIR. The purpose of this operation is to facilitate warpgroup-level matrix multiply and accumulate (WGMMA) operations on Hopper GPUs with sm_90a architecture. Previously, the `nvvm.wgmma.mma_async` operation was introduced to support warpgroup-level matrix operations in NVVM dialect. This op is used multiple instances of `nvvm.wgmma.mma_async` to achieve the desired shape. The new `nvgpu.warpgroup.mma` operation abstracts this complexity and provides a higher-level interface for performing warpgroup-level matrix operations. The `nvgpu.warpgroup.mma` does followings: 1) Corresponds multiple `wgmma` instructions. 2) Iterates input matrix descriptors to achieve the desired computation shape. 3) Groups and runs `wgmma` instructions asynchronously, and eventually waits them. This are done by `wgmma.fence.aligned`, `wgmma.commit.group.sync.aligned`, and `wgmma.wait.group.sync.aligned` 4) Results fragmented matrices Here's an example usage of the `nvgpu.warpgroup.mma` operation: ``` %wgmmaResult, %wgmmaResult2 = nvgpu.warpgroup.mma %descA, %descB, %acc1, %acc2 {transposeB}: !nvgpu.wgmma.descriptor<tensor = memref<128x64xf16, 3>>, !nvgpu.wgmma.descriptor<tensor = memref<64x128xf16, 3>>, !nvgpu.warpgroup.accumulator< fragmented = vector<64x128xf32>>, !nvgpu.warpgroup.accumulator< fragmented = vector<64x128xf32>> -> !nvgpu.warpgroup.accumulator< fragmented = vector<64x128xf32>>, !nvgpu.warpgroup.accumulator< fragmented = vector<64x128xf32>> ``` The op will result following PTX: ``` wgmma.fence.sync.aligned; wgmma.mma_async.sync.aligned.m64n128k16.f32.f16.f16 {%f1, %f2, 62 more registers}, %descA, %descB, p, 1, 1, 0, 1; wgmma.mma_async.sync.aligned.m64n128k16.f32.f16.f16 {%f1, %f2, 62 more registers}, %descA+2, %descB+128, p, 1, 1, 0, 1; wgmma.mma_async.sync.aligned.m64n128k16.f32.f16.f16 {%f1, %f2, 62 more registers}, %descA+4, %descB+256, p, 1, 1, 0, 1; wgmma.mma_async.sync.aligned.m64n128k16.f32.f16.f16 {%f1, %f2, 62 more registers}, %descA+8, %descB+348, p, 1, 1, 0, 1; wgmma.mma_async.sync.aligned.m64n128k16.f32.f16.f16 {%f500,%f501, 62 more registers}, %descA+512, %descB, p, 1, 1, 0, 1; wgmma.mma_async.sync.aligned.m64n128k16.f32.f16.f16 {%f500,%f501, 62 more registers}, %descA+514, %descB+128, p, 1, 1, 0, 1; wgmma.mma_async.sync.aligned.m64n128k16.f32.f16.f16 {%f500,%f501, 62 more registers}, %descA+516, %descB+256, p, 1, 1, 0, 1; wgmma.mma_async.sync.aligned.m64n128k16.f32.f16.f16 {%f500,%f501, 62 more registers}, %descA+518, %descB+348, p, 1, 1, 0, 1; wgmma.commit_group.sync.aligned; wgmma.wait_group.sync.aligned 1; ``` The Op keeps - first 64 registers (`{%f1, %f2, 62 more registers}`) -> `%acc1` - second 64 registers (`{%f500,%f501, 62 more registers}`) -> `%acc2`. -
Simon Pilgrim authored
[AArch64] Don't rely on (zext (trunc x)) pattern to detect zext_inreg MULL patterns - use value tracking directly As explained on D159533, I'm trying to generalize the "(zext (trunc x)) -> x iff the upper bits are known zero" fold in getNode() and I was seeing assertions in the aarch64 mull matching code as it was assuming these 'zero-extend-inreg' patterns will remain from earlier in LowerMUL. Instead I've updated selectUmullSmull/skipExtensionForVectorMULL to just use value tracking to detect when the upper bits are known zero, and to insert the truncation nodes later if necessary. Differential Revision: https://reviews.llvm.org/D159537
-
Simon Pilgrim authored
As discussed on D159537, using the SDValue operands directly instead of peeking inside to the SDNode prevents any issues where a non-zero result index has been used.
-
Matthew Devereau authored
rshrnb's shift amount operand must be between 1-EltSizeInBits. This patch stops RSHRNB ISD nodes being emitted in this case
-
Sam McCall authored
My immediate use for this is not in checked-in code, but rather the ability to plug printed flow conditions (from analysis logs) back into sat solver unittests to reproduce slowness. It does allow simplifying some of the existing solver tests, though.
-
Timm Bäder authored
-
Sam McCall authored
We're traversing the same chains of module ancestors and include locations repeatedly, despite already populating sets that can detect it! This is a problem because translateFile() is expensive. I think we can avoid it entirely, but this seems like an improvement either way. I removed a callback indirection rather than giving it a more complicated signature, and accordingly renamed the lambdas to be more concrete.
-
jeanPerier authored
Previous code was finalizing polymorphic components according to static type (calling the static type final routine, if any). There is no way (I think) to know from a Fortran::runtime::typeInfo::Component if an allocatable component is polymorphic or not. So this patch just always uses the dynamic type descriptor to check for derived type allocatable component finalization.
-
Ivan Kosarev authored
Reviewed By: rampitec, Joe_Nash Differential Revision: https://reviews.llvm.org/D156099
-
Benjamin Maxwell authored
Add support for following vector to tile (MOVA) intrinsics to ArmSME dialect: ``` llvm.aarch64.sme.read.vert llvm.aarch64.sme.read.horiz ``` This also slightly updates ArmSME_IntrOp to support return values.
-
Sam McCall authored
(In addition to the un-normalized form, so this is back-compatible)
-
Tobias Hieta authored
Doxygen documentation takes very long to build, when making releases we want to get the normal documentation up earlier so that we don't have to wait for doxygen documentation. This PR just adds the flag to disable doxygen builds, I will then later make a PR that changes the actions to first build the normal docs and another job to build the doxygen docs.
-
Guray Ozen authored
The test code had a space between "CHECK" and ":" that prohibits testing. This PR fixes this problem
-
Tobias Hieta authored
I landed this format helper, but unfortunately, it didn't work because of permissions, it could not add comments on a fork's PR. @cor3ntin informed me there are fixes for this that you had worked on @tstellar - but I didn't have time to read up on it too much. Can you explain what changes are needed to get the action to be able to write comments on fork's PR?
-
Nikita Popov authored
For pointer types, checking the address space is the same as type equality now, so we no longer need the special case.
-
David Spickett authored
This reverts commit 9a99944d. Due to test suite failures on all our SVE buildbots e.g.: https://lab.llvm.org/buildbot/#/builders/184/builds/7375 clang: ../llvm/llvm/lib/Target/AArch64/AArch64TargetTransformInfo.cpp:3565: InstructionCost llvm::AArch64TTIImpl::getShuffleCost(TTI::ShuffleKind, VectorType *, ArrayRef<int>, TTI::TargetCostKind, int, VectorType *, ArrayRef<const Value *>): Assertion `Mask.size() == TpNumElts && "Expected Mask and Tp size to match!"' failed.
-
Kazu Hirata authored
-
Balint Cristian authored
Enables summary descriptions along with the names of the feature. Descriptions here are simply looked up via the available llvm tablegen data.
-
Kazu Hirata authored
-
Nikita Popov authored
Remove all the expandCodeFor() uses that specify an explicit type, as well as InsertNoopCastOfTo() calls and most uses of getEffectiveSCEVType(). The only place where no-op casts can now be inserted are public expandCodeFor() uses.
-