summaryrefslogtreecommitdiff
path: root/llvm/lib/Target/AMDGPU/AMDGPUISelLowering.cpp
AgeCommit message (Collapse)Author
2025-06-26[AMDGPU] Convert 64-bit sra to 32-bit if shift amt >= 32 (#144421)LU-JOHN
Use KnownBits to convert 64-bit sra to 32-bit sra. Scaled-down alive2 verification with 16/8-bit types: https://alive2.llvm.org/ce/z/LamASk --------- Signed-off-by: John Lu <John.Lu@amd.com>
2025-06-23AMDGPU: Use reportFatalUsageError for unhandled calling conventions (#145261)Matt Arsenault
Should switch this to DiagnosticInfo and use the default calling convention, but that would require passing in the context.
2025-06-20AMDGPU: Remove AMDGPUInstrInfo class (#144984)Matt Arsenault
This was never constructed and only provided one static helper function.
2025-06-13[AMDGPU] Convert more 64-bit lshr to 32-bit if shift amt>=32 (#138204)LU-JOHN
Convert vector 64-bit lshr to 32-bit if shift amt is known to be >= 32. Also convert scalar 64-bit lshr to 32-bit if shift amt is variable but known to be >=32. --------- Signed-off-by: John Lu <John.Lu@amd.com>
2025-05-30[AMDGPU] Extend SRA i64 simplification for shift amts in range [33:62] (#138913)LU-JOHN
Extend sra i64 simplification to shift constants in range [33:62]. Shift amounts 32 and 63 were already handled. New testing for shift amts 33 and 62 added in sra.ll. Changes to other test files were to adapt previous test results to this extension. --------- Signed-off-by: John Lu <John.Lu@amd.com>
2025-05-28Warn on misuse of DiagnosticInfo classes that hold Twines (#137397)Justin Bogner
This annotates the `Twine` passed to the constructors of the various DiagnosticInfo subclasses with `[[clang::lifetimebound]]`, which causes us to warn when we would try to print the twine after it had already been destructed. We also update `DiagnosticInfoUnsupported` to hold a `const Twine &` like all of the other DiagnosticInfo classes, since this warning allows us to clean up all of the places where it was being used incorrectly.
2025-05-16[AMDGPU] Handle min/max in isNarrowingProfitable (#140206)Pierre van Houtryve
Introduces a slight regression in some cases but it'll even out once we disable the promotion in CGP.
2025-05-09AMDGPU: Use poison instead of undef in more lowerings (#139208)Matt Arsenault
2025-05-09AMDGPU: Handle minimumnum/maximumnum in fneg combines (#139133)Matt Arsenault
2025-05-08AMDGPU][True16][CodeGen] FP_Round f64 to f16 in true16 (#128911)Brox Chen
Update the f64 to f16 lowering for targets which support f16 types. For unsafe mode, lowered to two FP_ROUND. (This patch https://reviews.llvm.org/D154528 stops from combining these two FP_ROUND back). In safe mode, select LowerF64ToF16 (round-to-nearest-even rounding mode)
2025-05-05AMDGPU: Fix gcc -Wenum-compare warning (#138529)Matt Arsenault
2025-04-23[DAG] shouldReduceLoadWidth - add optional<unsigned> byte offset argument ↵Simon Pilgrim
(#136723) Based off feedback for #129695 - we need to be able to determine the load offset of smaller loads when trying to determine whether a multiple use load should be split (in particular for AVX subvector extractions). This patch adds a std::optional<unsigned> ByteOffset argument to shouldReduceLoadWidth calls for where we know the constant offset to allow targets to make use of it in future patches.
2025-04-18[DAG] isKnownNeverNaN - add DemandedElts element mask to isKnownNeverNaN ↵Simon Pilgrim
calls (#135952) Matches what we've done for computeKnownBits etc. to improve vector handling
2025-04-02[Clang][AMDGPU] Add __builtin_amdgcn_cvt_off_f32_i4 (#133741)Juan Manuel Martinez Caamaño
This built-in maps to `V_CVT_OFF_F32_I4` which treats its input as a 4-bit signed integer and returns `0.0625f * src`. SWDEV-518861
2025-03-29[GlobalISel][NFC] Rename GISelKnownBits to GISelValueTracking (#133466)Tim Gymnich
- rename `GISelKnownBits` to `GISelValueTracking` to analyze more than just `KnownBits` in the future
2025-03-28AMDGPU: Convert vector 64-bit shl to 32-bit if shift amt >= 32 (#132964)LU-JOHN
Convert vector 64-bit shl to 32-bit if shift amt is known to be >= 32. --------- Signed-off-by: John Lu <John.Lu@amd.com>
2025-03-20[AMDGPU] Dynamic VGPR support for llvm.amdgcn.cs.chain (#130094)Diana Picus
The llvm.amdgcn.cs.chain intrinsic has a 'flags' operand which may indicate that we want to reallocate the VGPRs before performing the call. A call with the following arguments: ``` llvm.amdgcn.cs.chain %callee, %exec, %sgpr_args, %vgpr_args, /*flags*/0x1, %num_vgprs, %fallback_exec, %fallback_callee ``` is supposed to do the following: - copy the SGPR and VGPR args into their respective registers - try to change the VGPR allocation - if the allocation has succeeded, set EXEC to %exec and jump to %callee, otherwise set EXEC to %fallback_exec and jump to %fallback_callee This patch implements the dynamic VGPR behaviour by generating an S_ALLOC_VGPR followed by S_CSELECT_B32/64 instructions for the EXEC and callee. The rest of the call sequence is left undisturbed (i.e. identical to the case where the flags are 0 and we don't use dynamic VGPRs). We achieve this by introducing some new pseudos (SI_CS_CHAIN_TC_Wn_DVGPR) which are expanded in the SILateBranchLowering pass, just like the simpler SI_CS_CHAIN_TC_Wn pseudos. The main reason is so that we don't risk other passes (particularly the PostRA scheduler) introducing instructions between the S_ALLOC_VGPR and the jump. Such instructions might end up using VGPRs that have been deallocated, or the wrong EXEC mask. Once the whole backend treats S_ALLOC_VGPR and changes to EXEC as barriers for instructions that use VGPRs, we could in principle move the expansion earlier (but in the absence of a good reason for that my personal preference is to keep it later in order to make debugging easier). Since the expansion happens after register allocation, we're careful to select constants to immediate operands instead of letting ISel generate S_MOVs which could interfere with register allocation (i.e. make it look like we need more registers than we actually do). For GFX12, S_ALLOC_VGPR only works in wave32 mode, so we bail out during ISel in wave64 mode. However, we can define the pseudos for wave64 too so it's easy to handle if future generations support it. --------- Co-authored-by: Ana Mihajlovic <Ana.Mihajlovic@amd.com> Co-authored-by: Matt Arsenault <Matthew.Arsenault@amd.com>
2025-03-18[CodeGen][GlobalISel] Add a getVectorIdxWidth and getVectorIdxLLT. (#131526)David Green
From #106446, this adds a variant of getVectorIdxTy that returns an LLT. Many uses only look at the width, so a getVectorIdxWidth was added as the common base.
2025-03-03[AMDGPU] Simplify conditional expressions. NFC. (#129228)Jay Foad
Simplfy `cond ? val : false` to `cond && val` and similar.
2025-02-13AMDGPU: Reduce shl64 to shl32 if shift range is [63-32] (#125574)LU-JOHN
Reduce: DST = shl i64 X, Y where Y is in the range [63-32] to: DST = [0, shl i32 X, (Y & 32)] Alive2 analysis: https://alive2.llvm.org/ce/z/w_u5je --------- Signed-off-by: John Lu <John.Lu@amd.com>
2025-02-04AMDGPU: Generalize truncate of shift of cast build_vector combine (#125617)Matt Arsenault
Previously we only handled cases that looked like the high element extract of a 64-bit shift. Generalize this to handle any multiple indexing. I was hoping this would help avoid some regressions, but it did not. It does however reduce the number of steps the DAG takes to process these cases. NFC-ish, I have yet to find an example where this changes the final output.
2024-12-21[SelectionDAG] Virtualize isTargetStrictFPOpcode / isTargetMemoryOpcode ↵Sergei Barannikov
(#119969) With this change, targets are no longer required to put memory / strict-fp opcodes after special `ISD::FIRST_TARGET_MEMORY_OPCODE`/`ISD::FIRST_TARGET_STRICTFP_OPCODE` markers. This will also allow autogenerating `isTargetMemoryOpcode`/`isTargetStrictFPOpcode (#119709). Pull Request: https://github.com/llvm/llvm-project/pull/119969
2024-12-18[SelectionDAG] Add SDNode::user_begin() and use it in some places (#120509)Craig Topper
Most of these are just places that want the first user and aren't iterating over the whole list. While there I changed some use_size() == 1 to hasOneUse() which is more efficient. This is part of an effort to rename use_iterator to user_iterator and provide a use_iterator that dereferences to SDUse&. This patch helps reduce the diff on later patches.
2024-12-18[SelectionDAG] Rename SDNode::uses() to users(). (#120499)Craig Topper
This function is most often used in range based loops or algorithms where the iterator is implicitly dereferenced. The dereference returns an SDNode * of the user rather than SDUse * so users() is a better name. I've long beeen annoyed that we can't write a range based loop over SDUse when we need getOperandNo. I plan to rename use_iterator to user_iterator and add a use_iterator that returns SDUse& on dereference. This will make it more like IR.
2024-12-10[DAGCombine] Remove oneuse restrictions for RISCV in folding (shl (add_nsw ↵LiqinWeng
x, c1)), c2) and folding (shl(sext(add x, c1)), c2) in some scenarios (#101294) This patch remove the restriction for folding (shl (add_nsw x, c1)), c2) and folding (shl(sext(add x, c1)), c2), and test case from dhrystone , see this link: riscv32: https://godbolt.org/z/o8GdMKrae riscv64: https://godbolt.org/z/Yh5bPz56z
2024-12-08Revert "[amdgpu][lds] Simplify error diag path - lds variable names are no ↵Jon Chesterfield
longer special" Test case didn't run locally, investigating This reverts commit 7bad469182ff2f6423ea209d5a1e81acca600568.
2024-12-08[amdgpu][lds] Simplify error diag path - lds variable names are no longer ↵Jon Chesterfield
special
2024-11-25[AMDGPU] Use getSignedConstant() where necessary (#117328)Nikita Popov
Create signed constant using getSignedConstant(), to avoid future assertion failures when we disable implicit truncation in getConstant(). This also touches some generic legalization code, which apparently only AMDGPU tests.
2024-11-13[AMDGPU] Remove unused includes (NFC) (#116154)Kazu Hirata
Identified with misc-include-cleaner.
2024-11-11[AMDGPU] Remove unused AMDGPUISD enum members (NFC) (#115582)Sergei Barannikov
Those were only used in `getTargetNodeName`.
2024-11-06[AMDGPU] modify named barrier builtins and intrinsics (#114550)Gang Chen
Use a local pointer type to represent the named barrier in builtin and intrinsic. This makes the definitions more user friendly bacause they do not need to worry about the hardware ID assignment. Also this approach is more like the other popular GPU programming language. Named barriers should be represented as global variables of addrspace(3) in LLVM-IR. Compiler assigns the special LDS offsets for those variables during AMDGPULowerModuleLDS pass. Those addresses are converted to hw barrier ID during instruction selection. The rest of the instruction-selection changes are primarily due to the intrinsic-definition changes.
2024-10-29DAG: Fix legalization of vector addrspacecasts (#113964)Matt Arsenault
2024-10-09[TTI] NFC: Port TLI.shouldSinkOperands to TTI (#110564)Jeffrey Byrnes
Porting to TTI provides direct access to the instruction cost model, which can enable instruction cost based sinking without introducing code duplication.
2024-10-03[AMDGPU] Qualify auto. NFC. (#110878)Jay Foad
Generated automatically with: $ clang-tidy -fix -checks=-*,llvm-qualified-auto $(find lib/Target/AMDGPU/ -type f)
2024-09-26[AMDGPU] Implement isSDNodeAlwaysUniform for INTRINSIC_W_CHAIN (#110114)Jay Foad
There are no always uniform side-effecting intrinsics upstream to test this with, but we have examples downstream.
2024-09-19[AMDGPU] Promote uniform ops to I32 in DAGISel (#106383)Pierre van Houtryve
Promote uniform binops, selects and setcc between 2 and 16 bits to 32 bits in DAGISel Solves #64591
2024-09-06[AMDGPU] Add target intrinsic for s_buffer_prefetch_data (#107293)Stanislav Mekhanoshin
2024-08-29AMDGPU: Use pattern to select instruction for intrinsic llvm.fptrunc.round ↵Changpeng Fang
(#105761) Use GCNPat instead of Custom Lowering to select instructions for intrinsic llvm.fptrunc.round. "SupportedRoundMode : TImmLeaf" is used as a predicate to select only when the rounding mode is supported. "as_hw_round_mode : SDNodeXForm" is developed to translate the round modes to the corresponding ones that hardware recognizes.
2024-08-29DAG: Check if is_fpclass is custom, instead of isLegalOrCustom (#105577)Matt Arsenault
For some reason, isOperationLegalOrCustom is not the same as isOperationLegal || isOperationCustom. Unfortunately, it checks if the type is legal which makes it uesless for custom lowering on non-legal types (which is always ppcf128). Really the DAG builder shouldn't be going to expand this in the builder, it makes it difficult to work with. It's only here to work around the DAG requiring legal integer types the same size as the FP type after type legalization.
2024-08-27[AMDGPU] Fix sign confusion in performMulLoHiCombine (#105831)Jay Foad
SMUL_LOHI and UMUL_LOHI are different operations because the high part of the result is different, so it is not OK to optimize the signed version to MUL_U24/MULHI_U24 or the unsigned version to MUL_I24/MULHI_I24.
2024-08-17AMDGPU: Add tonearest and towardzero roundings for intrinsic ↵Changpeng Fang
llvm.fptrunc.round (#104486) This work simplifies and generalizes the instruction definition for intrinsic llvm.fptrunc.round. We no longer name the instruction with the rounding mode. Instead, we introduce an immediate operand for the rounding mode for the pseudo instruction. This immediate will be used to set up the hardware mode register at the time the real instruction is generated. We name the pseudo instruction as FPTRUNC_ROUND_F16_F32 (for f32 -> f16), which is easy to generalize for other types. "round.towardzero" and "round.tonearest" are added for f32 -> f16 truncating, in addition to the existing "round.upward" and "round.downward". Other rounding modes are not supported by hardware at this moment.
2024-08-13[SelectionDAG] Replace EVTToAPFloatSemantics with MVT/EVT::getFltSemantics. ↵Craig Topper
(#103001)
2024-08-09[llvm] Construct SmallVector<SDValue> with ArrayRef (NFC) (#102578)Kazu Hirata
2024-08-08AMDGPU: Directly handle all atomicrmw cases in SIISelLowering (#102439)Matt Arsenault
2024-08-01AMDGPU: Cleanup extract_subvector actions (NFC) (#101454)Matt Arsenault
The base AMDGPUISelLowering was setting custom action on 16-bit vector types, but also set in SIISelLowering.
2024-07-25AMDGPU: Fix asserting in DAG kernel argument lowering on v6i32 (#100528)Matt Arsenault
Remove this pointless assertion for the number of vector elements.
2024-07-24[AMDGPU] Implement llvm.lrint intrinsic lowering (#98931)Sumanth Gundapaneni
This patch enabled the target-independent lowering of llvm.lrint via GlobalISel. For SelectionDAG, the instrinsic is custom lowered for AMDGPU.
2024-07-23[AMDGPU] Implement llvm.lround intrinsic lowering. (#98970)Sumanth Gundapaneni
This patch enables the target-independent lowering of llvm.lround via GlobalISel. For SelectionDAG, the instrinsic is custom lowered for AMDGPU. In order to support vector floating point input for llvm.lround, this patch extends the target independent APIs and provide support for scalarizing. pr98950 is needed to let verifier allow vector floating point types
2024-07-17[AMDGPU] Use range-based for loops. NFC. (#99047)Jay Foad
2024-07-16[AMDGPU] Use std::min with initializer list. NFC.Jay Foad