1
0
mirror of https://github.com/RPCS3/llvm-mirror.git synced 2024-11-26 12:43:36 +01:00
llvm-mirror/lib/Target/AMDGPU
Mirko Brkusanin 2afa995d4a [AMDGPU][GlobalISel] Legalize G_ABS
Legalize and select G_ABS so that we can use llvm.abs intrinsic

Differential Revision: https://reviews.llvm.org/D102391
2021-06-04 14:46:43 +02:00
..
AsmParser [AMDGPU][MC][NFC] Fixed typos in parser 2021-06-04 15:40:42 +03:00
Disassembler [AMDGPU] Add support for architected flat scratch 2021-05-14 10:53:48 -07:00
MCTargetDesc [AMDGPU] Add support for architected flat scratch 2021-05-14 10:53:48 -07:00
TargetInfo
Utils [AMDGPU] Fix kernel LDS lowering for constants 2021-05-26 11:34:50 -07:00
AMDGPU.h [AMDGPU] Experiments show that the GCNRegBankReassign pass significantly impacts 2021-04-26 17:21:49 -04:00
AMDGPU.td [AMDGPU] Add support for architected flat scratch 2021-05-14 10:53:48 -07:00
AMDGPUAliasAnalysis.cpp [NFC][AA] Prepare to convert AliasResult to class with PartialAlias offset. 2021-04-09 12:54:22 +03:00
AMDGPUAliasAnalysis.h
AMDGPUAlwaysInlinePass.cpp [AMDGPU] Disable forceful inline of non-kernel functions which use LDS. 2021-04-15 09:12:56 +05:30
AMDGPUAnnotateKernelFeatures.cpp [AMDGPU] Do not annotate features for graphics 2021-05-03 10:33:11 +02:00
AMDGPUAnnotateUniformValues.cpp [OpaquePtr] Clean up some uses of Type::getPointerElementType() 2021-05-31 09:54:57 -07:00
AMDGPUArgumentUsageInfo.cpp
AMDGPUArgumentUsageInfo.h [AMDGPU] gfx90a support 2021-02-17 16:01:32 -08:00
AMDGPUAsmPrinter.cpp [AMDGPU][NFC] Remove author's name from codebase 2021-06-02 00:51:48 +05:30
AMDGPUAsmPrinter.h AMDGPU: Add target id and code object v4 support 2021-03-24 11:54:05 -04:00
AMDGPUAtomicOptimizer.cpp [AMDGPU] Use reductions instead of scans in the atomic optimizer 2021-03-26 15:38:14 +00:00
AMDGPUCallingConv.td [AMDGPU] Fix large return values with amdgpu_gfx 2021-04-15 14:57:56 +02:00
AMDGPUCallLowering.cpp AMDGPU/GlobalISel: Use IncomingValueAssigner for implicit return 2021-05-27 11:28:52 -04:00
AMDGPUCallLowering.h AMDGPU/GlobalISel: Implement tail calls 2021-05-13 18:57:42 -04:00
AMDGPUCodeGenPrepare.cpp Normalize interaction with boolean attributes 2021-04-17 08:17:33 +02:00
AMDGPUCombine.td AMDGPU/GlobalISel: Remove redundant G_FCANONICALIZE 2021-04-27 12:26:37 +02:00
AMDGPUExportClustering.cpp [AMDGPU][MC] Refactored exp tgt handling 2021-01-26 14:54:15 +03:00
AMDGPUExportClustering.h [llvm] Add missing header guards (NFC) 2021-01-30 09:53:42 -08:00
AMDGPUFeatures.td
AMDGPUFixFunctionBitcasts.cpp
AMDGPUFrameLowering.cpp
AMDGPUFrameLowering.h
AMDGPUGenRegisterBankInfo.def
AMDGPUGISel.td AMDGPU/GlobalISel: Add integer med3 combines 2021-04-27 11:52:23 +02:00
AMDGPUGlobalISelUtils.cpp [AMDGPU][GlobalISel] Handle G_PTR_ADD when looking for constant offset 2021-01-28 11:20:09 +01:00
AMDGPUGlobalISelUtils.h [ADT] Move DenseMapInfo for ArrayRef/StringRef into respective headers (NFC) 2021-06-03 18:34:36 +02:00
AMDGPUHSAMetadataStreamer.cpp AMDGPU: Add target id and code object v4 support 2021-03-24 11:54:05 -04:00
AMDGPUHSAMetadataStreamer.h AMDGPU: Add target id and code object v4 support 2021-03-24 11:54:05 -04:00
AMDGPUInstCombineIntrinsic.cpp
AMDGPUInstrInfo.cpp
AMDGPUInstrInfo.h
AMDGPUInstrInfo.td [AMDGPU] Expose __builtin_amdgcn_perm for v_perm_b32 2021-05-06 16:17:33 -07:00
AMDGPUInstructions.td [AMDGPU] Improve Codegen for build_vector 2021-05-12 14:17:44 +01:00
AMDGPUInstructionSelector.cpp [AMDGPU] All GWS instructions need aligned VGPR on gfx90a 2021-06-01 17:08:03 -07:00
AMDGPUInstructionSelector.h [AMDGPU] Use SIInstrFlags for flat variants. NFC 2021-04-09 12:28:36 +02:00
AMDGPUISelDAGToDAG.cpp [AMDGPU] Fix function calls with flat scratch 2021-05-28 11:22:13 +02:00
AMDGPUISelLowering.cpp [AMDGPU] Fix a crash when selecting a particular case of buffer_load_format_d16 2021-06-03 16:40:18 -04:00
AMDGPUISelLowering.h [AMDGPU] Select V_CVT_*16_F16 more often 2021-05-05 08:57:51 +01:00
AMDGPULateCodeGenPrepare.cpp AMDGPU: Fix assert on constant load from addrspacecasted pointer 2021-05-11 20:12:20 -04:00
AMDGPULegalizerInfo.cpp [AMDGPU][GlobalISel] Legalize G_ABS 2021-06-04 14:46:43 +02:00
AMDGPULegalizerInfo.h [AMDGPU] Remove dead declaration (NFC). 2021-05-25 16:04:04 +05:30
AMDGPULibCalls.cpp Normalize interaction with boolean attributes 2021-04-17 08:17:33 +02:00
AMDGPULibFunc.cpp
AMDGPULibFunc.h
AMDGPULowerIntrinsics.cpp
AMDGPULowerKernelArguments.cpp
AMDGPULowerKernelAttributes.cpp Normalize interaction with boolean attributes 2021-04-17 08:17:33 +02:00
AMDGPULowerModuleLDSPass.cpp Revert "[AMDGPU] Increase alignment of LDS globals if necessary before LDS lowering." 2021-06-04 11:16:46 +05:30
AMDGPUMachineCFGStructurizer.cpp [AMDGPU] Fix build breakage 2021-02-14 09:02:55 -08:00
AMDGPUMachineFunction.cpp [AMDGPU] Fix module LDS selection 2021-05-20 15:59:01 -07:00
AMDGPUMachineFunction.h [amdgpu] Implement lower function LDS pass 2021-03-15 15:24:01 +00:00
AMDGPUMachineModuleInfo.cpp
AMDGPUMachineModuleInfo.h
AMDGPUMacroFusion.cpp
AMDGPUMacroFusion.h [llvm] Add missing header guards (NFC) 2021-01-30 09:53:42 -08:00
AMDGPUMCInstLower.cpp [AMDGPU] Remove SI_MASK_BRANCH 2021-03-09 09:13:23 +08:00
AMDGPUMIRFormatter.cpp [AMDGPU] Implement mir parseCustomPseudoSourceValue 2021-01-22 11:24:08 +01:00
AMDGPUMIRFormatter.h [AMDGPU] Implement mir parseCustomPseudoSourceValue 2021-01-22 11:24:08 +01:00
AMDGPUOpenCLEnqueuedBlockLowering.cpp
AMDGPUPerfHintAnalysis.cpp
AMDGPUPerfHintAnalysis.h
AMDGPUPostLegalizerCombiner.cpp AMDGPU/GlobalISel: Remove redundant G_FCANONICALIZE 2021-04-27 12:26:37 +02:00
AMDGPUPreLegalizerCombiner.cpp AMDGPU/GlobalISel: Add integer med3 combines 2021-04-27 11:52:23 +02:00
AMDGPUPrintfRuntimeBinding.cpp AMDGPUPrintfRuntimeBinding - don't dereference a dyn_cast<> pointer. NFCI. 2021-01-28 12:38:44 +00:00
AMDGPUPromoteAlloca.cpp [OpaquePtr] Create API to make a copy of a PointerType with some address space 2021-06-01 16:52:32 -07:00
AMDGPUPropagateAttributes.cpp [AMDGPU] [IndirectCalls] Don't propagate attributes to address taken functions and their callees 2021-06-04 11:36:56 +05:30
AMDGPUPTNote.h AMDGPU: Add target id and code object v4 support 2021-03-24 11:54:05 -04:00
AMDGPURegBankCombiner.cpp AMDGPU/GlobalISel: Add integer med3 combines 2021-04-27 11:52:23 +02:00
AMDGPURegisterBankInfo.cpp [AMDGPU][GlobalISel] Legalize G_ABS 2021-06-04 14:46:43 +02:00
AMDGPURegisterBankInfo.h [AMDGPU][GlobalISel] Use scalar min/max instructions 2021-02-04 17:04:32 +00:00
AMDGPURegisterBanks.td
AMDGPURewriteOutArguments.cpp
AMDGPUSearchableTables.td [AMDGPU] gfx90a support 2021-02-17 16:01:32 -08:00
AMDGPUSubtarget.cpp [AMDGPU] Add support for architected flat scratch 2021-05-14 10:53:48 -07:00
AMDGPUSubtarget.h [AMDGPU] gfx90a support 2021-02-17 16:01:32 -08:00
AMDGPUTargetMachine.cpp [AMDGPU] Disable the SIFormMemoryClauses pass at -O1 2021-05-12 11:51:59 -04:00
AMDGPUTargetMachine.h [NewPM] Hide pass manager debug logging behind -debug-pass-manager-verbose 2021-05-07 21:51:47 -07:00
AMDGPUTargetObjectFile.cpp
AMDGPUTargetObjectFile.h
AMDGPUTargetTransformInfo.cpp [TTI] NFC: Change getTypeLegalizationCost to return InstructionCost. 2021-04-30 22:51:51 +03:00
AMDGPUTargetTransformInfo.h [TTI] NFC: Change getVectorSplitCost to return InstructionCost 2021-04-21 17:32:02 +03:00
AMDGPUUnifyDivergentExitNodes.cpp [AMDGPU] Fix typo in comment 2021-05-18 10:15:49 +01:00
AMDGPUUnifyMetadata.cpp
AMDILCFGStructurizer.cpp [NewPM] Cleanup IR printing instrumentation 2021-04-15 09:50:55 -07:00
AMDKernelCodeT.h
BUFInstructions.td [AMDGPU] Fix extra waitcnt being added with BUFFER_INVL2 2021-05-11 13:17:33 -07:00
CaymanInstructions.td
CMakeLists.txt [AMDGPU] Experiments show that the GCNRegBankReassign pass significantly impacts 2021-04-26 17:21:49 -04:00
DSInstructions.td [AMDGPU] Re-arrange ds_read/ds_write ISel pattern for better readability. 2021-04-20 16:17:15 +05:30
EvergreenInstructions.td [AMDGPU] Inline FSHRPattern into its only use. NFC. 2021-03-26 09:32:02 +00:00
EXPInstructions.td
FLATInstructions.td AMDGPU: Restore atomic fp feature on FP atomic instruction definitions 2021-04-22 21:32:01 -04:00
GCNDPPCombine.cpp [AMDGPU] GCNDPPCombine: don't shrink V_ADD_CO_U32 if carry out is used 2021-04-20 09:17:52 +01:00
GCNHazardRecognizer.cpp [AMDGPU] Simplify getWaitStatesSince. NFC. 2021-04-30 08:58:24 +01:00
GCNHazardRecognizer.h [AMDGPU][NFC] Refactor hazard recognition IsHazardFn and IsExpiredFn 2021-04-30 09:18:56 +09:00
GCNILPSched.cpp
GCNIterativeScheduler.cpp
GCNIterativeScheduler.h
GCNMinRegStrategy.cpp
GCNNSAReassign.cpp [AMDGPU] Do not reassign spilled registers 2021-01-27 16:29:05 -08:00
GCNProcessors.td [AMDGPU] Add gfx1034 target 2021-05-13 14:25:18 -04:00
GCNRegPressure.cpp [AMDGPU] gfx90a support 2021-02-17 16:01:32 -08:00
GCNRegPressure.h [AMDGPU] gfx90a support 2021-02-17 16:01:32 -08:00
GCNSchedStrategy.cpp [AMDGPU] Avoid second rescheduling for some regions 2021-02-26 12:29:37 -08:00
GCNSchedStrategy.h [AMDGPU] Avoid second rescheduling for some regions 2021-02-26 12:29:37 -08:00
GCNSubtarget.h [AMDGPU] Add support for architected flat scratch 2021-05-14 10:53:48 -07:00
InstCombineTables.td
MIMGInstructions.td [AMDGPU][AsmParser/Disassembler] Correct A16 and G16 handling 2021-05-14 09:25:44 +01:00
R600.td
R600AsmPrinter.cpp
R600AsmPrinter.h
R600ClauseMergePass.cpp
R600ControlFlowFinalizer.cpp
R600Defines.h
R600EmitClauseMarkers.cpp
R600ExpandSpecialInstrs.cpp
R600FrameLowering.cpp
R600FrameLowering.h
R600InstrFormats.td
R600InstrInfo.cpp
R600InstrInfo.h
R600Instructions.td
R600ISelLowering.cpp [DAG] Add a generic expansion for SHIFT_PARTS opcodes using funnel shifts 2021-05-07 13:12:30 +01:00
R600ISelLowering.h [DAG] Add a generic expansion for SHIFT_PARTS opcodes using funnel shifts 2021-05-07 13:12:30 +01:00
R600MachineFunctionInfo.cpp
R600MachineFunctionInfo.h
R600MachineScheduler.cpp
R600MachineScheduler.h
R600OpenCLImageTypeLoweringPass.cpp TransformUtils: Fix metadata handling in CloneModule (and improve CloneFunctionInto) 2021-02-15 11:56:00 -08:00
R600OptimizeVectorRegisters.cpp
R600Packetizer.cpp
R600Processors.td
R600RegisterInfo.cpp
R600RegisterInfo.h
R600RegisterInfo.td
R600Schedule.td
R600Subtarget.h
R700Instructions.td
SIAnnotateControlFlow.cpp [AMDGPU] Do not annotate an else branch if there is a kill 2021-03-12 11:52:08 +09:00
SIDefines.h [AMDGPU] IsFlatScratch/Global -> FlatScratch/Global 2021-04-09 11:20:31 +02:00
SIFixSGPRCopies.cpp [AMDGPU] Introduce Strict WQM mode 2021-03-03 14:19:16 +01:00
SIFixVGPRCopies.cpp
SIFoldOperands.cpp [AMDGPU] SIFoldOperands: clean up tryConstantFoldOp 2021-05-06 09:55:22 +01:00
SIFormMemoryClauses.cpp [AMDGPU][NFC] Fix typos in SIFormMemoryClauses description 2021-05-06 07:47:39 -07:00
SIFrameLowering.cpp [WebAssembly][CodeGen] IR support for WebAssembly local variables 2021-06-01 11:31:39 +02:00
SIFrameLowering.h
SIInsertHardClauses.cpp [AMDGPU] Do not clause NSA instructions 2021-05-14 12:54:56 +09:00
SIInsertWaitcnts.cpp [AMDGPU] Remove assert 2021-05-12 14:52:37 +02:00
SIInstrFormats.td [AMDGPU] IsFlatScratch/Global -> FlatScratch/Global 2021-04-09 11:20:31 +02:00
SIInstrInfo.cpp [AMDGPU] All GWS instructions need aligned VGPR on gfx90a 2021-06-01 17:08:03 -07:00
SIInstrInfo.h [AMDGPU] Update SCC defs to VCC when uses are changed to VCC 2021-05-14 18:05:05 -04:00
SIInstrInfo.td [AMDGPU] More accurate names for dpp operand types 2021-05-25 10:35:25 -04:00
SIInstructions.td [AMDGPU] Improve Codegen for build_vector 2021-05-12 14:17:44 +01:00
SIISelLowering.cpp [AMDGPU] All GWS instructions need aligned VGPR on gfx90a 2021-06-01 17:08:03 -07:00
SIISelLowering.h [TTI] NFC: Change getTypeLegalizationCost to return InstructionCost. 2021-04-30 22:51:51 +03:00
SILateBranchLowering.cpp [AMDGPU] Add MDT update missing from D98915 2021-03-20 13:38:58 +09:00
SILoadStoreOptimizer.cpp AMDGPU: Fix SILoadStoreOptimizer for gfx90a 2021-05-11 21:26:43 -04:00
SILowerControlFlow.cpp [AMDGPU] Add llvm.amdgcn.wqm.demote intrinsic 2021-02-15 08:45:46 +09:00
SILowerI1Copies.cpp [Target] Use llvm::append_range (NFC) 2021-01-24 12:18:56 -08:00
SILowerSGPRSpills.cpp [AMDGPU] Free reserved VGPR if no SGPR spill 2021-03-12 08:11:14 +08:00
SIMachineFunctionInfo.cpp [AMDGPU] Add support for architected flat scratch 2021-05-14 10:53:48 -07:00
SIMachineFunctionInfo.h [AMDGPU] Serialize MFInfo::ScavengeFI 2021-05-07 11:15:25 +02:00
SIMachineScheduler.cpp
SIMachineScheduler.h
SIMemoryLegalizer.cpp [AMDGPU] Update gfx90a memory model support 2021-04-07 22:17:58 +00:00
SIModeRegister.cpp
SIOptimizeExecMasking.cpp [AMDGPU] Move kill lowering to WQM pass and add live mask tracking 2021-02-11 20:31:29 +09:00
SIOptimizeExecMaskingPreRA.cpp [AMDGPU] SIOptimizeExecMaskingPreRA should check constant bus constraint when folds EXEC copy 2021-03-24 14:14:13 +03:00
SIPeepholeSDWA.cpp
SIPostRABundler.cpp AMDGPU: Use kill instruction to hint soft clause live ranges 2021-02-26 18:26:40 -05:00
SIPreAllocateWWMRegs.cpp [AMDGPU] Save WWM registers in functions 2021-04-23 18:09:24 +02:00
SIPreEmitPeephole.cpp [AMDGPU] Remove set_gpr_idx instructions in conditional blocks 2021-04-30 22:15:45 +01:00
SIProgramInfo.cpp
SIProgramInfo.h [AMDGPU] gfx90a support 2021-02-17 16:01:32 -08:00
SIRegisterInfo.cpp [AMDGPU] Allow buildSpillLoadStore in empty bb 2021-04-29 12:53:20 +02:00
SIRegisterInfo.h [AMDGPU][NFC] Remove non-existing function header 2021-05-26 18:20:33 +02:00
SIRegisterInfo.td [AMDGPU] Use if instead of foreach in a few places. NFC. 2021-04-20 14:20:30 +01:00
SISchedule.td [AMDGPU] Add TransVALU to gfx10 2021-04-20 15:34:43 +02:00
SIShrinkInstructions.cpp [AMDGPU] Fix v_swap_b32 formation on physical registers 2021-04-29 20:53:40 +01:00
SIWholeQuadMode.cpp [AMDGPU] Fix WQM failure with single block inactive demote 2021-05-06 21:02:26 +09:00
SMInstructions.td [AMDGPU][MC][GFX9] Corrected SMEM decoding 2021-04-06 14:10:46 +03:00
SOPInstructions.td [AMDGPU][GlobalISel] Legalize G_ABS 2021-06-04 14:46:43 +02:00
VIInstrFormats.td
VOP1Instructions.td [AMDGPU] Remove redundant field from DPP8 def 2021-04-16 16:23:52 -04:00
VOP2Instructions.td [AMDGPU] Remove redundant field from DPP8 def 2021-04-16 16:23:52 -04:00
VOP3Instructions.td [AMDGPU] Tweak VOP3_INTERP16 profile 2021-05-17 15:28:00 +01:00
VOP3PInstructions.td [AMDGPU] Refactor VOP3P Profile and AsmParser, NFC 2021-04-16 13:06:50 -04:00
VOPCInstructions.td [AMDGPU] Add new EmitDstSel field to VOPPofile. NFC. 2021-04-15 12:07:08 -07:00
VOPInstructions.td [AMDGPU] Set unused dst_sel to '?' in the encoding 2021-05-17 08:38:52 -07:00