aboutsummaryrefslogtreecommitdiffstats
diff options
context:
space:
mode:
authorDimitry Andric <dim@FreeBSD.org>2017-06-10 13:44:06 +0000
committerDimitry Andric <dim@FreeBSD.org>2017-06-10 13:44:06 +0000
commit7ab83427af0f77b59941ceba41d509d7d097b065 (patch)
treecc41c05b1db454e3d802f34df75e636ee922ad87
parentd288ef4c1788d3a951a7558c68312c2d320612b1 (diff)
downloadsrc-7ab83427af0f77b59941ceba41d509d7d097b065.tar.gz
src-7ab83427af0f77b59941ceba41d509d7d097b065.zip
Vendor import of llvm trunk r305145:vendor/llvm/llvm-trunk-r305145
Notes
Notes: svn path=/vendor/llvm/dist/; revision=319780 svn path=/vendor/llvm/llvm-trunk-r305145/; revision=319781; tag=vendor/llvm/llvm-trunk-r305145
-rw-r--r--CMakeLists.txt6
-rw-r--r--bindings/go/llvm/ir.go6
-rw-r--r--bindings/go/llvm/ir_test.go26
-rw-r--r--bindings/ocaml/llvm/llvm.ml2
-rw-r--r--bindings/ocaml/llvm/llvm.mli3
-rw-r--r--bindings/ocaml/llvm/llvm_ocaml.c14
-rw-r--r--cmake/modules/HandleLLVMOptions.cmake7
-rw-r--r--cmake/modules/TableGen.cmake50
-rw-r--r--docs/AMDGPUUsage.rst3566
-rw-r--r--docs/CodeGenerator.rst59
-rw-r--r--docs/CompilerWriterInfo.rst11
-rw-r--r--docs/LangRef.rst33
-rw-r--r--docs/ReleaseNotes.rst11
-rw-r--r--docs/index.rst4
-rw-r--r--examples/ExceptionDemo/ExceptionDemo.cpp4
-rw-r--r--include/llvm-c/Core.h14
-rw-r--r--include/llvm-c/ExecutionEngine.h2
-rw-r--r--include/llvm-c/Support.h2
-rw-r--r--include/llvm-c/TargetMachine.h2
-rw-r--r--include/llvm/ADT/APInt.h5
-rw-r--r--include/llvm/ADT/GraphTraits.h1
-rw-r--r--include/llvm/ADT/ImmutableSet.h6
-rw-r--r--include/llvm/ADT/PointerUnion.h2
-rw-r--r--include/llvm/ADT/PostOrderIterator.h2
-rw-r--r--include/llvm/ADT/PriorityWorklist.h2
-rw-r--r--include/llvm/ADT/SCCIterator.h10
-rw-r--r--include/llvm/ADT/SmallPtrSet.h6
-rw-r--r--include/llvm/ADT/SmallVector.h21
-rw-r--r--include/llvm/ADT/SparseMultiSet.h2
-rw-r--r--include/llvm/ADT/StringExtras.h2
-rw-r--r--include/llvm/ADT/StringRef.h2
-rw-r--r--include/llvm/ADT/iterator_range.h2
-rw-r--r--include/llvm/Analysis/AliasAnalysis.h4
-rw-r--r--include/llvm/Analysis/AssumptionCache.h2
-rw-r--r--include/llvm/Analysis/BranchProbabilityInfo.h11
-rw-r--r--include/llvm/Analysis/ConstantFolding.h6
-rw-r--r--include/llvm/Analysis/DemandedBits.h2
-rw-r--r--include/llvm/Analysis/InlineCost.h2
-rw-r--r--include/llvm/Analysis/InstructionSimplify.h6
-rw-r--r--include/llvm/Analysis/LazyBranchProbabilityInfo.h9
-rw-r--r--include/llvm/Analysis/LazyValueInfo.h7
-rw-r--r--include/llvm/Analysis/LoopInfoImpl.h2
-rw-r--r--include/llvm/Analysis/MemoryDependenceAnalysis.h2
-rw-r--r--include/llvm/Analysis/MemorySSAUpdater.h2
-rw-r--r--include/llvm/Analysis/ObjCARCAnalysisUtils.h2
-rw-r--r--include/llvm/Analysis/ObjCARCInstKind.h2
-rw-r--r--include/llvm/Analysis/ScalarEvolutionNormalization.h2
-rw-r--r--include/llvm/Analysis/TargetLibraryInfo.h8
-rw-r--r--include/llvm/Analysis/TargetTransformInfo.h31
-rw-r--r--include/llvm/Analysis/TargetTransformInfoImpl.h18
-rw-r--r--include/llvm/BinaryFormat/COFF.h713
-rw-r--r--include/llvm/BinaryFormat/Dwarf.def (renamed from include/llvm/Support/Dwarf.def)0
-rw-r--r--include/llvm/BinaryFormat/Dwarf.h (renamed from include/llvm/Support/Dwarf.h)58
-rw-r--r--include/llvm/BinaryFormat/ELF.h (renamed from include/llvm/Support/ELF.h)23
-rw-r--r--include/llvm/BinaryFormat/ELFRelocs/AArch64.def (renamed from include/llvm/Support/ELFRelocs/AArch64.def)0
-rw-r--r--include/llvm/BinaryFormat/ELFRelocs/AMDGPU.def (renamed from include/llvm/Support/ELFRelocs/AMDGPU.def)0
-rw-r--r--include/llvm/BinaryFormat/ELFRelocs/ARM.def (renamed from include/llvm/Support/ELFRelocs/ARM.def)0
-rw-r--r--include/llvm/BinaryFormat/ELFRelocs/AVR.def (renamed from include/llvm/Support/ELFRelocs/AVR.def)0
-rw-r--r--include/llvm/BinaryFormat/ELFRelocs/BPF.def (renamed from include/llvm/Support/ELFRelocs/BPF.def)0
-rw-r--r--include/llvm/BinaryFormat/ELFRelocs/Hexagon.def (renamed from include/llvm/Support/ELFRelocs/Hexagon.def)0
-rw-r--r--include/llvm/BinaryFormat/ELFRelocs/Lanai.def (renamed from include/llvm/Support/ELFRelocs/Lanai.def)0
-rw-r--r--include/llvm/BinaryFormat/ELFRelocs/Mips.def (renamed from include/llvm/Support/ELFRelocs/Mips.def)0
-rw-r--r--include/llvm/BinaryFormat/ELFRelocs/PowerPC.def (renamed from include/llvm/Support/ELFRelocs/PowerPC.def)0
-rw-r--r--include/llvm/BinaryFormat/ELFRelocs/PowerPC64.def (renamed from include/llvm/Support/ELFRelocs/PowerPC64.def)0
-rw-r--r--include/llvm/BinaryFormat/ELFRelocs/RISCV.def (renamed from include/llvm/Support/ELFRelocs/RISCV.def)0
-rw-r--r--include/llvm/BinaryFormat/ELFRelocs/Sparc.def (renamed from include/llvm/Support/ELFRelocs/Sparc.def)0
-rw-r--r--include/llvm/BinaryFormat/ELFRelocs/SystemZ.def (renamed from include/llvm/Support/ELFRelocs/SystemZ.def)0
-rw-r--r--include/llvm/BinaryFormat/ELFRelocs/WebAssembly.def (renamed from include/llvm/Support/ELFRelocs/WebAssembly.def)0
-rw-r--r--include/llvm/BinaryFormat/ELFRelocs/i386.def (renamed from include/llvm/Support/ELFRelocs/i386.def)0
-rw-r--r--include/llvm/BinaryFormat/ELFRelocs/x86_64.def (renamed from include/llvm/Support/ELFRelocs/x86_64.def)0
-rw-r--r--include/llvm/BinaryFormat/MachO.def (renamed from include/llvm/Support/MachO.def)0
-rw-r--r--include/llvm/BinaryFormat/MachO.h1984
-rw-r--r--include/llvm/BinaryFormat/Magic.h73
-rw-r--r--include/llvm/BinaryFormat/Wasm.h (renamed from include/llvm/Support/Wasm.h)46
-rw-r--r--include/llvm/BinaryFormat/WasmRelocs/WebAssembly.def (renamed from include/llvm/Support/WasmRelocs/WebAssembly.def)0
-rw-r--r--include/llvm/Bitcode/BitcodeReader.h13
-rw-r--r--include/llvm/Bitcode/LLVMBitCodes.h2
-rw-r--r--include/llvm/CodeGen/BasicTTIImpl.h76
-rw-r--r--include/llvm/CodeGen/DFAPacketizer.h37
-rw-r--r--include/llvm/CodeGen/DIE.h2
-rw-r--r--include/llvm/CodeGen/ExecutionDepsFix.h42
-rw-r--r--include/llvm/CodeGen/FastISel.h39
-rw-r--r--include/llvm/CodeGen/FunctionLoweringInfo.h33
-rw-r--r--include/llvm/CodeGen/GCMetadata.h38
-rw-r--r--include/llvm/CodeGen/GCMetadataPrinter.h25
-rw-r--r--include/llvm/CodeGen/GCStrategy.h2
-rw-r--r--include/llvm/CodeGen/GlobalISel/InstructionSelector.h2
-rw-r--r--include/llvm/CodeGen/GlobalISel/LegalizerHelper.h2
-rw-r--r--include/llvm/CodeGen/GlobalISel/MachineIRBuilder.h2
-rw-r--r--include/llvm/CodeGen/LexicalScopes.h3
-rw-r--r--include/llvm/CodeGen/LiveInterval.h4
-rw-r--r--include/llvm/CodeGen/LiveRegUnits.h2
-rw-r--r--include/llvm/CodeGen/MIRParser/MIRParser.h20
-rw-r--r--include/llvm/CodeGen/MIRYamlMapping.h140
-rw-r--r--include/llvm/CodeGen/MachineBasicBlock.h4
-rw-r--r--include/llvm/CodeGen/MachineFunction.h4
-rw-r--r--include/llvm/CodeGen/MachineFunctionInitializer.h38
-rw-r--r--include/llvm/CodeGen/MachineFunctionPass.h2
-rw-r--r--include/llvm/CodeGen/MachineMemOperand.h2
-rw-r--r--include/llvm/CodeGen/MachineModuleInfo.h15
-rw-r--r--include/llvm/CodeGen/MachineModuleInfoImpls.h4
-rw-r--r--include/llvm/CodeGen/MachineOperand.h2
-rw-r--r--include/llvm/CodeGen/MachineRegisterInfo.h4
-rw-r--r--include/llvm/CodeGen/RegAllocRegistry.h17
-rw-r--r--include/llvm/CodeGen/RegisterPressure.h10
-rw-r--r--include/llvm/CodeGen/RegisterUsageInfo.h17
-rw-r--r--include/llvm/CodeGen/ScheduleDAG.h2
-rw-r--r--include/llvm/CodeGen/ScheduleDAGInstrs.h2
-rw-r--r--include/llvm/CodeGen/ScheduleDFS.h2
-rw-r--r--include/llvm/CodeGen/SchedulerRegistry.h17
-rw-r--r--include/llvm/CodeGen/SelectionDAG.h39
-rw-r--r--include/llvm/CodeGen/SelectionDAGNodes.h4
-rw-r--r--include/llvm/CodeGen/SlotIndexes.h2
-rw-r--r--include/llvm/CodeGen/StackProtector.h19
-rw-r--r--include/llvm/CodeGen/TailDuplicator.h37
-rw-r--r--include/llvm/CodeGen/TargetLoweringObjectFileImpl.h8
-rw-r--r--include/llvm/CodeGen/TargetPassConfig.h11
-rw-r--r--include/llvm/CodeGen/TargetSchedule.h3
-rw-r--r--include/llvm/CodeGen/VirtRegMap.h9
-rw-r--r--include/llvm/Config/abi-breaking.h.cmake3
-rw-r--r--include/llvm/Config/config.h.cmake21
-rw-r--r--include/llvm/DebugInfo/CodeView/CVRecord.h6
-rw-r--r--include/llvm/DebugInfo/CodeView/CodeView.h18
-rw-r--r--include/llvm/DebugInfo/CodeView/DebugChecksumsSubsection.h4
-rw-r--r--include/llvm/DebugInfo/CodeView/DebugCrossExSubsection.h64
-rw-r--r--include/llvm/DebugInfo/CodeView/DebugCrossImpSubsection.h88
-rw-r--r--include/llvm/DebugInfo/CodeView/DebugInlineeLinesSubsection.h7
-rw-r--r--include/llvm/DebugInfo/CodeView/DebugLinesSubsection.h10
-rw-r--r--include/llvm/DebugInfo/CodeView/DebugStringTableSubsection.h3
-rw-r--r--include/llvm/DebugInfo/CodeView/DebugSubsectionRecord.h14
-rw-r--r--include/llvm/DebugInfo/CodeView/DebugSubsectionVisitor.h128
-rw-r--r--include/llvm/DebugInfo/CodeView/DebugSymbolRVASubsection.h59
-rw-r--r--include/llvm/DebugInfo/CodeView/DebugSymbolsSubsection.h3
-rw-r--r--include/llvm/DebugInfo/CodeView/EnumTables.h2
-rw-r--r--include/llvm/DebugInfo/CodeView/TypeSerializer.h2
-rw-r--r--include/llvm/DebugInfo/CodeView/TypeTableBuilder.h5
-rw-r--r--include/llvm/DebugInfo/DIContext.h2
-rw-r--r--include/llvm/DebugInfo/DWARF/DWARFAbbreviationDeclaration.h4
-rw-r--r--include/llvm/DebugInfo/DWARF/DWARFAcceleratorTable.h2
-rw-r--r--include/llvm/DebugInfo/DWARF/DWARFAttribute.h2
-rw-r--r--include/llvm/DebugInfo/DWARF/DWARFCompileUnit.h6
-rw-r--r--include/llvm/DebugInfo/DWARF/DWARFContext.h35
-rw-r--r--include/llvm/DebugInfo/DWARF/DWARFDebugInfoEntry.h2
-rw-r--r--include/llvm/DebugInfo/DWARF/DWARFDebugPubTable.h2
-rw-r--r--include/llvm/DebugInfo/DWARF/DWARFDebugRangeList.h2
-rw-r--r--include/llvm/DebugInfo/DWARF/DWARFDie.h7
-rw-r--r--include/llvm/DebugInfo/DWARF/DWARFFormValue.h2
-rw-r--r--include/llvm/DebugInfo/DWARF/DWARFTypeUnit.h2
-rw-r--r--include/llvm/DebugInfo/DWARF/DWARFUnit.h32
-rw-r--r--include/llvm/DebugInfo/MSF/MappedBlockStream.h1
-rw-r--r--include/llvm/DebugInfo/PDB/DIA/DIAEnumDebugStreams.h1
-rw-r--r--include/llvm/DebugInfo/PDB/DIA/DIAEnumLineNumbers.h1
-rw-r--r--include/llvm/DebugInfo/PDB/DIA/DIAEnumSourceFiles.h1
-rw-r--r--include/llvm/DebugInfo/PDB/DIA/DIAEnumSymbols.h1
-rw-r--r--include/llvm/DebugInfo/PDB/Native/DbiModuleDescriptor.h5
-rw-r--r--include/llvm/DebugInfo/PDB/Native/DbiStream.h3
-rw-r--r--include/llvm/DebugInfo/PDB/PDBSymbol.h4
-rw-r--r--include/llvm/ExecutionEngine/ExecutionEngine.h2
-rw-r--r--include/llvm/ExecutionEngine/Orc/CompileOnDemandLayer.h2
-rw-r--r--include/llvm/ExecutionEngine/Orc/ExecutionUtils.h2
-rw-r--r--include/llvm/ExecutionEngine/Orc/IRCompileLayer.h2
-rw-r--r--include/llvm/ExecutionEngine/Orc/RTDyldObjectLinkingLayer.h2
-rw-r--r--include/llvm/ExecutionEngine/RTDyldMemoryManager.h2
-rw-r--r--include/llvm/IR/Attributes.h4
-rw-r--r--include/llvm/IR/BasicBlock.h4
-rw-r--r--include/llvm/IR/CallSite.h4
-rw-r--r--include/llvm/IR/Constants.h6
-rw-r--r--include/llvm/IR/DataLayout.h2
-rw-r--r--include/llvm/IR/DebugInfoMetadata.h11
-rw-r--r--include/llvm/IR/DiagnosticInfo.h4
-rw-r--r--include/llvm/IR/Dominators.h6
-rw-r--r--include/llvm/IR/Function.h4
-rw-r--r--include/llvm/IR/GetElementPtrTypeIterator.h4
-rw-r--r--include/llvm/IR/GlobalValue.h2
-rw-r--r--include/llvm/IR/GlobalVariable.h9
-rw-r--r--include/llvm/IR/IRBuilder.h26
-rw-r--r--include/llvm/IR/InstrTypes.h2
-rw-r--r--include/llvm/IR/Instruction.h8
-rw-r--r--include/llvm/IR/Instructions.h101
-rw-r--r--include/llvm/IR/IntrinsicsAMDGPU.td10
-rw-r--r--include/llvm/IR/Metadata.h8
-rw-r--r--include/llvm/IR/Module.h4
-rw-r--r--include/llvm/IR/ModuleSummaryIndex.h4
-rw-r--r--include/llvm/IR/OperandTraits.h3
-rw-r--r--include/llvm/IR/PatternMatch.h2
-rw-r--r--include/llvm/IR/Statepoint.h2
-rw-r--r--include/llvm/IR/Type.h2
-rw-r--r--include/llvm/IR/Use.h2
-rw-r--r--include/llvm/IR/Value.h2
-rw-r--r--include/llvm/LTO/LTO.h9
-rw-r--r--include/llvm/LinkAllIR.h2
-rw-r--r--include/llvm/LinkAllPasses.h4
-rw-r--r--include/llvm/MC/MCAsmInfo.h8
-rw-r--r--include/llvm/MC/MCAssembler.h6
-rw-r--r--include/llvm/MC/MCCodeView.h4
-rw-r--r--include/llvm/MC/MCContext.h2
-rw-r--r--include/llvm/MC/MCELFObjectWriter.h2
-rw-r--r--include/llvm/MC/MCFragment.h2
-rw-r--r--include/llvm/MC/MCMachObjectWriter.h4
-rw-r--r--include/llvm/MC/MCObjectFileInfo.h4
-rw-r--r--include/llvm/MC/MCParser/MCAsmParser.h2
-rw-r--r--include/llvm/MC/MCSection.h2
-rw-r--r--include/llvm/MC/MCSectionMachO.h2
-rw-r--r--include/llvm/MC/MCSymbolWasm.h2
-rw-r--r--include/llvm/MC/MCTargetOptions.h6
-rw-r--r--include/llvm/MC/MCWasmObjectWriter.h30
-rw-r--r--include/llvm/Object/Archive.h2
-rw-r--r--include/llvm/Object/COFF.h9
-rw-r--r--include/llvm/Object/COFFModuleDefinition.h2
-rw-r--r--include/llvm/Object/Decompressor.h2
-rw-r--r--include/llvm/Object/ELF.h2
-rw-r--r--include/llvm/Object/ELFObjectFile.h2
-rw-r--r--include/llvm/Object/ELFTypes.h2
-rw-r--r--include/llvm/Object/IRObjectFile.h15
-rw-r--r--include/llvm/Object/IRSymtab.h15
-rw-r--r--include/llvm/Object/MachO.h6
-rw-r--r--include/llvm/Object/MachOUniversal.h2
-rw-r--r--include/llvm/Object/ObjectFile.h10
-rw-r--r--include/llvm/Object/RelocVisitor.h4
-rw-r--r--include/llvm/Object/SymbolicFile.h7
-rw-r--r--include/llvm/Object/Wasm.h2
-rw-r--r--include/llvm/Object/WindowsResource.h74
-rw-r--r--include/llvm/ObjectYAML/COFFYAML.h2
-rw-r--r--include/llvm/ObjectYAML/CodeViewYAMLDebugSections.h32
-rw-r--r--include/llvm/ObjectYAML/CodeViewYAMLTypes.h5
-rw-r--r--include/llvm/ObjectYAML/DWARFYAML.h14
-rw-r--r--include/llvm/ObjectYAML/ELFYAML.h2
-rw-r--r--include/llvm/ObjectYAML/MachOYAML.h8
-rw-r--r--include/llvm/ObjectYAML/WasmYAML.h2
-rw-r--r--include/llvm/Option/ArgList.h2
-rw-r--r--include/llvm/Pass.h2
-rw-r--r--include/llvm/ProfileData/Coverage/CoverageMapping.h6
-rw-r--r--include/llvm/Support/AMDGPUCodeObjectMetadata.h (renamed from lib/Target/AMDGPU/MCTargetDesc/AMDGPUCodeObjectMetadata.h)6
-rw-r--r--include/llvm/Support/BinaryStreamArray.h258
-rw-r--r--include/llvm/Support/BinaryStreamReader.h20
-rw-r--r--include/llvm/Support/CBindingWrapping.h2
-rw-r--r--include/llvm/Support/COFF.h724
-rw-r--r--include/llvm/Support/Casting.h45
-rw-r--r--include/llvm/Support/CommandLine.h77
-rw-r--r--include/llvm/Support/ConvertUTF.h2
-rw-r--r--include/llvm/Support/DataTypes.h.cmake4
-rw-r--r--include/llvm/Support/Endian.h168
-rw-r--r--include/llvm/Support/Error.h29
-rw-r--r--include/llvm/Support/ErrorOr.h18
-rw-r--r--include/llvm/Support/FileSystem.h55
-rw-r--r--include/llvm/Support/FormatVariadic.h2
-rw-r--r--include/llvm/Support/GCOV.h4
-rw-r--r--include/llvm/Support/GenericDomTree.h2
-rw-r--r--include/llvm/Support/LowLevelTypeImpl.h2
-rw-r--r--include/llvm/Support/MachO.h2038
-rw-r--r--include/llvm/Support/MathExtras.h2
-rw-r--r--include/llvm/Support/MemoryBuffer.h4
-rw-r--r--include/llvm/Support/Solaris.h2
-rw-r--r--include/llvm/Support/SourceMgr.h2
-rw-r--r--include/llvm/Support/StringPool.h24
-rw-r--r--include/llvm/Support/TargetRegistry.h4
-rw-r--r--include/llvm/Support/raw_sha1_ostream.h4
-rw-r--r--include/llvm/Support/type_traits.h20
-rw-r--r--include/llvm/Target/TargetInstrInfo.h2
-rw-r--r--include/llvm/Target/TargetLowering.h69
-rw-r--r--include/llvm/Target/TargetMachine.h12
-rw-r--r--include/llvm/Target/TargetOptions.h10
-rw-r--r--include/llvm/Target/TargetSubtargetInfo.h2
-rw-r--r--include/llvm/Transforms/IPO/FunctionAttrs.h2
-rw-r--r--include/llvm/Transforms/Scalar/GVNExpression.h13
-rw-r--r--include/llvm/Transforms/Utils/EscapeEnumerator.h2
-rw-r--r--include/llvm/Transforms/Utils/FunctionComparator.h2
-rw-r--r--include/llvm/Transforms/Utils/ImportedFunctionsInliningStatistics.h2
-rw-r--r--include/llvm/Transforms/Utils/Local.h2
-rw-r--r--include/llvm/Transforms/Utils/LoopVersioning.h2
-rw-r--r--include/llvm/Transforms/Utils/OrderedInstructions.h51
-rw-r--r--include/llvm/Transforms/Utils/ValueMapper.h2
-rw-r--r--include/llvm/Transforms/Vectorize/SLPVectorizer.h2
-rw-r--r--include/llvm/module.modulemap44
-rw-r--r--lib/Analysis/AliasAnalysisEvaluator.cpp2
-rw-r--r--lib/Analysis/AliasSetTracker.cpp2
-rw-r--r--lib/Analysis/BranchProbabilityInfo.cpp47
-rw-r--r--lib/Analysis/CFLGraph.h1
-rw-r--r--lib/Analysis/CallPrinter.cpp2
-rw-r--r--lib/Analysis/CaptureTracking.cpp2
-rw-r--r--lib/Analysis/CodeMetrics.cpp2
-rw-r--r--lib/Analysis/ConstantFolding.cpp42
-rw-r--r--lib/Analysis/GlobalsModRef.cpp16
-rw-r--r--lib/Analysis/InlineCost.cpp4
-rw-r--r--lib/Analysis/InstCount.cpp3
-rw-r--r--lib/Analysis/InstructionSimplify.cpp68
-rw-r--r--lib/Analysis/LLVMBuild.txt2
-rw-r--r--lib/Analysis/LazyBranchProbabilityInfo.cpp8
-rw-r--r--lib/Analysis/LazyCallGraph.cpp3
-rw-r--r--lib/Analysis/LazyValueInfo.cpp194
-rw-r--r--lib/Analysis/Lint.cpp2
-rw-r--r--lib/Analysis/MemDepPrinter.cpp2
-rw-r--r--lib/Analysis/MemDerefPrinter.cpp4
-rw-r--r--lib/Analysis/MemoryDependenceAnalysis.cpp6
-rw-r--r--lib/Analysis/MemorySSAUpdater.cpp19
-rw-r--r--lib/Analysis/ModuleDebugInfoPrinter.cpp2
-rw-r--r--lib/Analysis/ModuleSummaryAnalysis.cpp10
-rw-r--r--lib/Analysis/ObjCARCInstKind.cpp2
-rw-r--r--lib/Analysis/RegionPrinter.cpp4
-rw-r--r--lib/Analysis/ScalarEvolution.cpp4
-rw-r--r--lib/Analysis/ScalarEvolutionNormalization.cpp2
-rw-r--r--lib/Analysis/TargetTransformInfo.cpp8
-rw-r--r--lib/Analysis/ValueTracking.cpp22
-rw-r--r--lib/Analysis/VectorUtils.cpp8
-rw-r--r--lib/AsmParser/LLParser.cpp4
-rw-r--r--lib/AsmParser/LLVMBuild.txt2
-rw-r--r--lib/BinaryFormat/CMakeLists.txt8
-rw-r--r--lib/BinaryFormat/Dwarf.cpp (renamed from lib/Support/Dwarf.cpp)200
-rw-r--r--lib/BinaryFormat/LLVMBuild.txt22
-rw-r--r--lib/BinaryFormat/Magic.cpp216
-rw-r--r--lib/Bitcode/Reader/BitcodeReader.cpp26
-rw-r--r--lib/Bitcode/Reader/MetadataLoader.cpp2
-rw-r--r--lib/Bitcode/Writer/BitcodeWriter.cpp10
-rw-r--r--lib/CMakeLists.txt1
-rw-r--r--lib/CodeGen/Analysis.cpp2
-rw-r--r--lib/CodeGen/AsmPrinter/ARMException.cpp2
-rw-r--r--lib/CodeGen/AsmPrinter/AsmPrinter.cpp10
-rw-r--r--lib/CodeGen/AsmPrinter/AsmPrinterDwarf.cpp2
-rw-r--r--lib/CodeGen/AsmPrinter/AsmPrinterInlineAsm.cpp2
-rw-r--r--lib/CodeGen/AsmPrinter/CodeViewDebug.cpp72
-rw-r--r--lib/CodeGen/AsmPrinter/CodeViewDebug.h51
-rw-r--r--lib/CodeGen/AsmPrinter/DIEHash.cpp4
-rw-r--r--lib/CodeGen/AsmPrinter/DebugLocStream.h2
-rw-r--r--lib/CodeGen/AsmPrinter/DwarfAccelTable.h2
-rw-r--r--lib/CodeGen/AsmPrinter/DwarfCFIException.cpp2
-rw-r--r--lib/CodeGen/AsmPrinter/DwarfCompileUnit.h2
-rw-r--r--lib/CodeGen/AsmPrinter/DwarfDebug.cpp2
-rw-r--r--lib/CodeGen/AsmPrinter/DwarfExpression.cpp2
-rw-r--r--lib/CodeGen/AsmPrinter/DwarfUnit.cpp4
-rw-r--r--lib/CodeGen/AsmPrinter/ErlangGCPrinter.cpp4
-rw-r--r--lib/CodeGen/AsmPrinter/LLVMBuild.txt2
-rw-r--r--lib/CodeGen/AsmPrinter/OcamlGCPrinter.cpp23
-rw-r--r--lib/CodeGen/AsmPrinter/WinException.cpp4
-rw-r--r--lib/CodeGen/BasicTargetTransformInfo.cpp2
-rw-r--r--lib/CodeGen/BranchFolding.cpp36
-rw-r--r--lib/CodeGen/BranchRelaxation.cpp6
-rw-r--r--lib/CodeGen/BuiltinGCs.cpp2
-rw-r--r--lib/CodeGen/CalcSpillWeights.cpp2
-rw-r--r--lib/CodeGen/CodeGen.cpp2
-rw-r--r--lib/CodeGen/CodeGenPrepare.cpp294
-rw-r--r--lib/CodeGen/DFAPacketizer.cpp83
-rw-r--r--lib/CodeGen/DeadMachineInstructionElim.cpp2
-rw-r--r--lib/CodeGen/DwarfEHPrepare.cpp4
-rw-r--r--lib/CodeGen/ExpandISelPseudos.cpp2
-rw-r--r--lib/CodeGen/ExpandPostRAPseudos.cpp2
-rw-r--r--lib/CodeGen/ExpandReductions.cpp6
-rw-r--r--lib/CodeGen/FaultMaps.cpp2
-rw-r--r--lib/CodeGen/FuncletLayout.cpp2
-rw-r--r--lib/CodeGen/GCMetadata.cpp16
-rw-r--r--lib/CodeGen/GCMetadataPrinter.cpp7
-rw-r--r--lib/CodeGen/GlobalISel/IRTranslator.cpp19
-rw-r--r--lib/CodeGen/GlobalISel/Legalizer.cpp1
-rw-r--r--lib/CodeGen/GlobalMerge.cpp3
-rw-r--r--lib/CodeGen/IfConversion.cpp2
-rw-r--r--lib/CodeGen/ImplicitNullChecks.cpp10
-rw-r--r--lib/CodeGen/InlineSpiller.cpp31
-rw-r--r--lib/CodeGen/LLVMTargetMachine.cpp141
-rw-r--r--lib/CodeGen/LexicalScopes.cpp2
-rw-r--r--lib/CodeGen/LiveIntervalAnalysis.cpp4
-rw-r--r--lib/CodeGen/LiveIntervalUnion.cpp4
-rw-r--r--lib/CodeGen/LiveRegMatrix.cpp6
-rw-r--r--lib/CodeGen/LocalStackSlotAllocation.cpp2
-rw-r--r--lib/CodeGen/MIRParser/MILexer.h2
-rw-r--r--lib/CodeGen/MIRParser/MIParser.cpp44
-rw-r--r--lib/CodeGen/MIRParser/MIParser.h16
-rw-r--r--lib/CodeGen/MIRParser/MIRParser.cpp138
-rw-r--r--lib/CodeGen/MIRPrinter.cpp65
-rw-r--r--lib/CodeGen/MIRPrintingPass.cpp4
-rw-r--r--lib/CodeGen/MachineBlockPlacement.cpp4
-rw-r--r--lib/CodeGen/MachineCSE.cpp2
-rw-r--r--lib/CodeGen/MachineCopyPropagation.cpp2
-rw-r--r--lib/CodeGen/MachineDominanceFrontier.cpp1
-rw-r--r--lib/CodeGen/MachineDominators.cpp2
-rw-r--r--lib/CodeGen/MachineFunction.cpp3
-rw-r--r--lib/CodeGen/MachineFunctionPass.cpp2
-rw-r--r--lib/CodeGen/MachineFunctionPrinterPass.cpp2
-rw-r--r--lib/CodeGen/MachineInstr.cpp6
-rw-r--r--lib/CodeGen/MachineLICM.cpp2
-rw-r--r--lib/CodeGen/MachineModuleInfo.cpp16
-rw-r--r--lib/CodeGen/MachineOutliner.cpp6
-rw-r--r--lib/CodeGen/MachinePipeliner.cpp2
-rw-r--r--lib/CodeGen/MachineRegionInfo.cpp2
-rw-r--r--lib/CodeGen/MachineRegisterInfo.cpp2
-rw-r--r--lib/CodeGen/MachineScheduler.cpp8
-rw-r--r--lib/CodeGen/MachineSink.cpp2
-rw-r--r--lib/CodeGen/MachineTraceMetrics.cpp2
-rw-r--r--lib/CodeGen/MachineVerifier.cpp11
-rw-r--r--lib/CodeGen/OptimizePHIs.cpp2
-rw-r--r--lib/CodeGen/PatchableFunction.cpp2
-rw-r--r--lib/CodeGen/PeepholeOptimizer.cpp2
-rw-r--r--lib/CodeGen/PostRAHazardRecognizer.cpp2
-rw-r--r--lib/CodeGen/RegAllocBase.cpp5
-rw-r--r--lib/CodeGen/RegAllocBasic.cpp2
-rw-r--r--lib/CodeGen/RegAllocGreedy.cpp96
-rw-r--r--lib/CodeGen/RegAllocPBQP.cpp16
-rw-r--r--lib/CodeGen/RegisterClassInfo.cpp2
-rw-r--r--lib/CodeGen/RegisterPressure.cpp4
-rw-r--r--lib/CodeGen/RegisterUsageInfo.cpp15
-rw-r--r--lib/CodeGen/RenameIndependentSubregs.cpp9
-rw-r--r--lib/CodeGen/ResetMachineFunctionPass.cpp2
-rw-r--r--lib/CodeGen/ScheduleDAG.cpp6
-rw-r--r--lib/CodeGen/ScheduleDAGInstrs.cpp6
-rw-r--r--lib/CodeGen/ScheduleDAGPrinter.cpp2
-rw-r--r--lib/CodeGen/SelectionDAG/DAGCombiner.cpp32
-rw-r--r--lib/CodeGen/SelectionDAG/FastISel.cpp2
-rw-r--r--lib/CodeGen/SelectionDAG/LegalizeDAG.cpp8
-rw-r--r--lib/CodeGen/SelectionDAG/ScheduleDAGFast.cpp2
-rw-r--r--lib/CodeGen/SelectionDAG/ScheduleDAGRRList.cpp2
-rw-r--r--lib/CodeGen/SelectionDAG/ScheduleDAGVLIW.cpp2
-rw-r--r--lib/CodeGen/SelectionDAG/SelectionDAG.cpp19
-rw-r--r--lib/CodeGen/SelectionDAG/SelectionDAGBuilder.cpp249
-rw-r--r--lib/CodeGen/SelectionDAG/SelectionDAGBuilder.h14
-rw-r--r--lib/CodeGen/SelectionDAG/SelectionDAGDumper.cpp2
-rw-r--r--lib/CodeGen/SelectionDAG/SelectionDAGISel.cpp28
-rw-r--r--lib/CodeGen/SelectionDAG/SelectionDAGPrinter.cpp2
-rw-r--r--lib/CodeGen/SelectionDAG/StatepointLowering.cpp4
-rw-r--r--lib/CodeGen/SelectionDAG/TargetLowering.cpp6
-rw-r--r--lib/CodeGen/ShadowStackGCLowering.cpp2
-rw-r--r--lib/CodeGen/SjLjEHPrepare.cpp2
-rw-r--r--lib/CodeGen/StackMaps.cpp2
-rw-r--r--lib/CodeGen/StackProtector.cpp8
-rw-r--r--lib/CodeGen/StackSlotColoring.cpp2
-rw-r--r--lib/CodeGen/TailDuplication.cpp14
-rw-r--r--lib/CodeGen/TailDuplicator.cpp38
-rw-r--r--lib/CodeGen/TargetFrameLoweringImpl.cpp14
-rw-r--r--lib/CodeGen/TargetLoweringBase.cpp8
-rw-r--r--lib/CodeGen/TargetLoweringObjectFileImpl.cpp122
-rw-r--r--lib/CodeGen/TargetOptionsImpl.cpp4
-rw-r--r--lib/CodeGen/TargetPassConfig.cpp71
-rw-r--r--lib/CodeGen/TargetRegisterInfo.cpp2
-rw-r--r--lib/CodeGen/TargetSchedule.cpp24
-rw-r--r--lib/CodeGen/TargetSubtargetInfo.cpp2
-rw-r--r--lib/CodeGen/VirtRegMap.cpp11
-rw-r--r--lib/CodeGen/WinEHPrepare.cpp2
-rw-r--r--lib/CodeGen/XRayInstrumentation.cpp30
-rw-r--r--lib/DebugInfo/CodeView/CMakeLists.txt11
-rw-r--r--lib/DebugInfo/CodeView/DebugChecksumsSubsection.cpp4
-rw-r--r--lib/DebugInfo/CodeView/DebugCrossExSubsection.cpp51
-rw-r--r--lib/DebugInfo/CodeView/DebugCrossImpSubsection.cpp91
-rw-r--r--lib/DebugInfo/CodeView/DebugInlineeLinesSubsection.cpp9
-rw-r--r--lib/DebugInfo/CodeView/DebugLinesSubsection.cpp11
-rw-r--r--lib/DebugInfo/CodeView/DebugStringTableSubsection.cpp14
-rw-r--r--lib/DebugInfo/CodeView/DebugSubsectionRecord.cpp28
-rw-r--r--lib/DebugInfo/CodeView/DebugSubsectionVisitor.cpp84
-rw-r--r--lib/DebugInfo/CodeView/DebugSymbolRVASubsection.cpp31
-rw-r--r--lib/DebugInfo/CodeView/TypeTableCollection.cpp3
-rw-r--r--lib/DebugInfo/DWARF/DWARFAbbreviationDeclaration.cpp5
-rw-r--r--lib/DebugInfo/DWARF/DWARFAcceleratorTable.cpp5
-rw-r--r--lib/DebugInfo/DWARF/DWARFCompileUnit.cpp4
-rw-r--r--lib/DebugInfo/DWARF/DWARFContext.cpp149
-rw-r--r--lib/DebugInfo/DWARF/DWARFDebugFrame.cpp8
-rw-r--r--lib/DebugInfo/DWARF/DWARFDebugInfoEntry.cpp2
-rw-r--r--lib/DebugInfo/DWARF/DWARFDebugLine.cpp2
-rw-r--r--lib/DebugInfo/DWARF/DWARFDebugLoc.cpp4
-rw-r--r--lib/DebugInfo/DWARF/DWARFDebugMacro.cpp4
-rw-r--r--lib/DebugInfo/DWARF/DWARFDebugPubTable.cpp4
-rw-r--r--lib/DebugInfo/DWARF/DWARFDebugRangeList.cpp2
-rw-r--r--lib/DebugInfo/DWARF/DWARFDie.cpp41
-rw-r--r--lib/DebugInfo/DWARF/DWARFFormValue.cpp10
-rw-r--r--lib/DebugInfo/DWARF/DWARFGdbIndex.cpp2
-rw-r--r--lib/DebugInfo/DWARF/DWARFTypeUnit.cpp2
-rw-r--r--lib/DebugInfo/DWARF/DWARFUnit.cpp60
-rw-r--r--lib/DebugInfo/DWARF/DWARFUnitIndex.cpp4
-rw-r--r--lib/DebugInfo/DWARF/LLVMBuild.txt2
-rw-r--r--lib/DebugInfo/PDB/DIA/DIAEnumDebugStreams.cpp4
-rw-r--r--lib/DebugInfo/PDB/DIA/DIAEnumLineNumbers.cpp2
-rw-r--r--lib/DebugInfo/PDB/DIA/DIAEnumSourceFiles.cpp2
-rw-r--r--lib/DebugInfo/PDB/DIA/DIAEnumSymbols.cpp2
-rw-r--r--lib/DebugInfo/PDB/DIA/DIARawSymbol.cpp7
-rw-r--r--lib/DebugInfo/PDB/Native/DbiModuleDescriptorBuilder.cpp16
-rw-r--r--lib/DebugInfo/PDB/Native/DbiStream.cpp6
-rw-r--r--lib/DebugInfo/PDB/Native/DbiStreamBuilder.cpp2
-rw-r--r--lib/DebugInfo/PDB/Native/ModuleDebugStream.cpp4
-rw-r--r--lib/DebugInfo/PDB/Native/PublicsStream.cpp10
-rw-r--r--lib/DebugInfo/PDB/PDBContext.cpp2
-rw-r--r--lib/DebugInfo/PDB/PDBSymbolBlock.cpp2
-rw-r--r--lib/DebugInfo/PDB/PDBSymbolCompilandDetails.cpp2
-rw-r--r--lib/DebugInfo/PDB/PDBSymbolCompilandEnv.cpp2
-rw-r--r--lib/DebugInfo/PDB/PDBSymbolCustom.cpp2
-rw-r--r--lib/DebugInfo/PDB/PDBSymbolFunc.cpp2
-rw-r--r--lib/DebugInfo/PDB/PDBSymbolFuncDebugEnd.cpp2
-rw-r--r--lib/DebugInfo/PDB/PDBSymbolFuncDebugStart.cpp2
-rw-r--r--lib/DebugInfo/PDB/PDBSymbolPublicSymbol.cpp2
-rw-r--r--lib/DebugInfo/PDB/PDBSymbolTypeBaseClass.cpp2
-rw-r--r--lib/DebugInfo/PDB/PDBSymbolTypeCustom.cpp2
-rw-r--r--lib/DebugInfo/PDB/PDBSymbolTypeDimension.cpp2
-rw-r--r--lib/DebugInfo/PDB/PDBSymbolTypeFriend.cpp2
-rw-r--r--lib/DebugInfo/PDB/PDBSymbolTypeFunctionSig.cpp2
-rw-r--r--lib/DebugInfo/PDB/PDBSymbolTypeManaged.cpp2
-rw-r--r--lib/DebugInfo/PDB/PDBSymbolTypeVTableShape.cpp2
-rw-r--r--lib/DebugInfo/PDB/PDBSymbolUnknown.cpp2
-rw-r--r--lib/DebugInfo/PDB/PDBSymbolUsingNamespace.cpp2
-rw-r--r--lib/DebugInfo/Symbolize/SymbolizableObjectFile.cpp2
-rw-r--r--lib/DebugInfo/Symbolize/Symbolize.cpp4
-rw-r--r--lib/ExecutionEngine/IntelJITEvents/IntelJITEventListener.cpp2
-rw-r--r--lib/ExecutionEngine/IntelJITEvents/jitprofiling.c2
-rw-r--r--lib/ExecutionEngine/Interpreter/ExternalFunctions.cpp4
-rw-r--r--lib/ExecutionEngine/OProfileJIT/OProfileJITEventListener.cpp2
-rw-r--r--lib/ExecutionEngine/Orc/IndirectionUtils.cpp2
-rw-r--r--lib/ExecutionEngine/Orc/OrcMCJITReplacement.h4
-rw-r--r--lib/ExecutionEngine/RuntimeDyld/RuntimeDyld.cpp4
-rw-r--r--lib/ExecutionEngine/RuntimeDyld/RuntimeDyldELF.cpp2
-rw-r--r--lib/ExecutionEngine/RuntimeDyld/RuntimeDyldImpl.h2
-rw-r--r--lib/ExecutionEngine/RuntimeDyld/Targets/RuntimeDyldCOFFI386.h4
-rw-r--r--lib/ExecutionEngine/RuntimeDyld/Targets/RuntimeDyldCOFFThumb.h4
-rw-r--r--lib/ExecutionEngine/RuntimeDyld/Targets/RuntimeDyldCOFFX86_64.h4
-rw-r--r--lib/ExecutionEngine/RuntimeDyld/Targets/RuntimeDyldELFMips.cpp2
-rw-r--r--lib/ExecutionEngine/SectionMemoryManager.cpp2
-rw-r--r--lib/Fuzzer/FuzzerDriver.cpp4
-rw-r--r--lib/Fuzzer/FuzzerExtFunctionsDlsymWin.cpp2
-rw-r--r--lib/Fuzzer/FuzzerLoop.cpp2
-rw-r--r--lib/Fuzzer/FuzzerMerge.cpp4
-rw-r--r--lib/Fuzzer/FuzzerMutate.cpp2
-rw-r--r--lib/Fuzzer/FuzzerShmemPosix.cpp6
-rw-r--r--lib/Fuzzer/FuzzerShmemWindows.cpp4
-rw-r--r--lib/Fuzzer/FuzzerTracePC.cpp2
-rw-r--r--lib/Fuzzer/FuzzerTraceState.cpp2
-rw-r--r--lib/Fuzzer/FuzzerUtilWindows.cpp2
-rw-r--r--lib/Fuzzer/afl/afl_driver.cpp12
-rw-r--r--lib/Fuzzer/test/AFLDriverTest.cpp2
-rw-r--r--lib/Fuzzer/test/AbsNegAndConstant64Test.cpp6
-rw-r--r--lib/Fuzzer/test/AbsNegAndConstantTest.cpp6
-rw-r--r--lib/Fuzzer/test/AccumulateAllocationsTest.cpp2
-rw-r--r--lib/Fuzzer/test/BadStrcmpTest.cpp2
-rw-r--r--lib/Fuzzer/test/BufferOverflowOnInput.cpp2
-rw-r--r--lib/Fuzzer/test/CallerCalleeTest.cpp2
-rw-r--r--lib/Fuzzer/test/CleanseTest.cpp2
-rw-r--r--lib/Fuzzer/test/CustomMutatorTest.cpp2
-rw-r--r--lib/Fuzzer/test/CxxStringEqTest.cpp4
-rw-r--r--lib/Fuzzer/test/DSOTestMain.cpp2
-rw-r--r--lib/Fuzzer/test/DivTest.cpp2
-rw-r--r--lib/Fuzzer/test/FourIndependentBranchesTest.cpp2
-rw-r--r--lib/Fuzzer/test/FullCoverageSetTest.cpp2
-rw-r--r--lib/Fuzzer/test/FuzzerUnittest.cpp4
-rw-r--r--lib/Fuzzer/test/LeakTest.cpp2
-rw-r--r--lib/Fuzzer/test/LeakTimeoutTest.cpp2
-rw-r--r--lib/Fuzzer/test/LoadTest.cpp2
-rw-r--r--lib/Fuzzer/test/Memcmp64BytesTest.cpp2
-rw-r--r--lib/Fuzzer/test/MemcmpTest.cpp2
-rw-r--r--lib/Fuzzer/test/NotinstrumentedTest.cpp2
-rw-r--r--lib/Fuzzer/test/NthRunCrashTest.cpp2
-rw-r--r--lib/Fuzzer/test/NullDerefOnEmptyTest.cpp2
-rw-r--r--lib/Fuzzer/test/NullDerefTest.cpp2
-rw-r--r--lib/Fuzzer/test/OneHugeAllocTest.cpp2
-rw-r--r--lib/Fuzzer/test/OutOfMemorySingleLargeMallocTest.cpp2
-rw-r--r--lib/Fuzzer/test/OutOfMemoryTest.cpp2
-rw-r--r--lib/Fuzzer/test/RepeatedBytesTest.cpp2
-rw-r--r--lib/Fuzzer/test/RepeatedMemcmp.cpp3
-rw-r--r--lib/Fuzzer/test/ShrinkControlFlowTest.cpp4
-rw-r--r--lib/Fuzzer/test/ShrinkValueProfileTest.cpp4
-rw-r--r--lib/Fuzzer/test/SignedIntOverflowTest.cpp4
-rw-r--r--lib/Fuzzer/test/SimpleCmpTest.cpp2
-rw-r--r--lib/Fuzzer/test/SimpleDictionaryTest.cpp2
-rw-r--r--lib/Fuzzer/test/SimpleHashTest.cpp2
-rw-r--r--lib/Fuzzer/test/SimpleTest.cpp2
-rw-r--r--lib/Fuzzer/test/SimpleThreadedTest.cpp2
-rw-r--r--lib/Fuzzer/test/SingleByteInputTest.cpp4
-rw-r--r--lib/Fuzzer/test/SingleMemcmpTest.cpp2
-rw-r--r--lib/Fuzzer/test/SingleStrcmpTest.cpp2
-rw-r--r--lib/Fuzzer/test/SingleStrncmpTest.cpp2
-rw-r--r--lib/Fuzzer/test/SpamyTest.cpp2
-rw-r--r--lib/Fuzzer/test/StrcmpTest.cpp4
-rw-r--r--lib/Fuzzer/test/StrncmpOOBTest.cpp4
-rw-r--r--lib/Fuzzer/test/StrncmpTest.cpp2
-rw-r--r--lib/Fuzzer/test/StrstrTest.cpp4
-rw-r--r--lib/Fuzzer/test/SwapCmpTest.cpp2
-rw-r--r--lib/Fuzzer/test/Switch2Test.cpp4
-rw-r--r--lib/Fuzzer/test/SwitchTest.cpp4
-rw-r--r--lib/Fuzzer/test/TableLookupTest.cpp4
-rw-r--r--lib/Fuzzer/test/ThreadedLeakTest.cpp2
-rw-r--r--lib/Fuzzer/test/ThreadedTest.cpp2
-rw-r--r--lib/Fuzzer/test/TimeoutEmptyTest.cpp2
-rw-r--r--lib/Fuzzer/test/TimeoutTest.cpp2
-rw-r--r--lib/Fuzzer/test/TraceMallocTest.cpp2
-rw-r--r--lib/Fuzzer/test/TwoDifferentBugsTest.cpp2
-rw-r--r--lib/IR/AsmWriter.cpp2
-rw-r--r--lib/IR/Attributes.cpp4
-rw-r--r--lib/IR/Comdat.cpp2
-rw-r--r--lib/IR/ConstantRange.cpp38
-rw-r--r--lib/IR/Constants.cpp13
-rw-r--r--lib/IR/Core.cpp12
-rw-r--r--lib/IR/DIBuilder.cpp4
-rw-r--r--lib/IR/DataLayout.cpp2
-rw-r--r--lib/IR/DebugInfo.cpp2
-rw-r--r--lib/IR/DebugLoc.cpp2
-rw-r--r--lib/IR/DiagnosticInfo.cpp6
-rw-r--r--lib/IR/DiagnosticPrinter.cpp4
-rw-r--r--lib/IR/Dominators.cpp20
-rw-r--r--lib/IR/Function.cpp4
-rw-r--r--lib/IR/Globals.cpp8
-rw-r--r--lib/IR/IRBuilder.cpp34
-rw-r--r--lib/IR/InlineAsm.cpp2
-rw-r--r--lib/IR/Instruction.cpp10
-rw-r--r--lib/IR/Instructions.cpp135
-rw-r--r--lib/IR/IntrinsicInst.cpp2
-rw-r--r--lib/IR/LLVMBuild.txt2
-rw-r--r--lib/IR/LLVMContext.cpp2
-rw-r--r--lib/IR/LLVMContextImpl.h2
-rw-r--r--lib/IR/LegacyPassManager.cpp2
-rw-r--r--lib/IR/Metadata.cpp4
-rw-r--r--lib/IR/Module.cpp6
-rw-r--r--lib/IR/OptBisect.cpp2
-rw-r--r--lib/IR/Type.cpp2
-rw-r--r--lib/IR/TypeFinder.cpp2
-rw-r--r--lib/IR/ValueSymbolTable.cpp2
-rw-r--r--lib/IR/Verifier.cpp6
-rw-r--r--lib/LLVMBuild.txt1
-rw-r--r--lib/LTO/LTO.cpp74
-rw-r--r--lib/LTO/ThinLTOCodeGenerator.cpp2
-rw-r--r--lib/MC/ELFObjectWriter.cpp20
-rw-r--r--lib/MC/MCAsmBackend.cpp2
-rw-r--r--lib/MC/MCAsmInfo.cpp2
-rw-r--r--lib/MC/MCAsmInfoDarwin.cpp2
-rw-r--r--lib/MC/MCAsmInfoELF.cpp2
-rw-r--r--lib/MC/MCAssembler.cpp4
-rw-r--r--lib/MC/MCCodeView.cpp2
-rw-r--r--lib/MC/MCContext.cpp8
-rw-r--r--lib/MC/MCDisassembler/Disassembler.cpp2
-rw-r--r--lib/MC/MCDisassembler/MCRelocationInfo.cpp2
-rw-r--r--lib/MC/MCDwarf.cpp6
-rw-r--r--lib/MC/MCELFStreamer.cpp6
-rw-r--r--lib/MC/MCExpr.cpp6
-rw-r--r--lib/MC/MCFragment.cpp4
-rw-r--r--lib/MC/MCInstPrinter.cpp2
-rw-r--r--lib/MC/MCInstrAnalysis.cpp2
-rw-r--r--lib/MC/MCMachOStreamer.cpp2
-rw-r--r--lib/MC/MCNullStreamer.cpp2
-rw-r--r--lib/MC/MCObjectFileInfo.cpp20
-rw-r--r--lib/MC/MCObjectWriter.cpp4
-rw-r--r--lib/MC/MCParser/AsmLexer.cpp2
-rw-r--r--lib/MC/MCParser/AsmParser.cpp8
-rw-r--r--lib/MC/MCParser/COFFAsmParser.cpp4
-rw-r--r--lib/MC/MCParser/DarwinAsmParser.cpp6
-rw-r--r--lib/MC/MCParser/ELFAsmParser.cpp4
-rw-r--r--lib/MC/MCParser/MCAsmLexer.cpp2
-rw-r--r--lib/MC/MCParser/MCAsmParser.cpp2
-rw-r--r--lib/MC/MCParser/MCTargetAsmParser.cpp2
-rw-r--r--lib/MC/MCRegisterInfo.cpp2
-rw-r--r--lib/MC/MCSection.cpp2
-rw-r--r--lib/MC/MCSectionCOFF.cpp2
-rw-r--r--lib/MC/MCSectionELF.cpp4
-rw-r--r--lib/MC/MCStreamer.cpp6
-rw-r--r--lib/MC/MCSubtargetInfo.cpp2
-rw-r--r--lib/MC/MCSymbol.cpp2
-rw-r--r--lib/MC/MCSymbolELF.cpp4
-rw-r--r--lib/MC/MCTargetOptions.cpp2
-rw-r--r--lib/MC/MCWasmObjectTargetWriter.cpp10
-rw-r--r--lib/MC/MCWinEH.cpp4
-rw-r--r--lib/MC/MachObjectWriter.cpp4
-rw-r--r--lib/MC/StringTableBuilder.cpp4
-rw-r--r--lib/MC/SubtargetFeature.cpp2
-rw-r--r--lib/MC/WasmObjectWriter.cpp267
-rw-r--r--lib/MC/WinCOFFObjectWriter.cpp6
-rw-r--r--lib/MC/WinCOFFStreamer.cpp4
-rw-r--r--lib/Object/Archive.cpp2
-rw-r--r--lib/Object/ArchiveWriter.cpp3
-rw-r--r--lib/Object/Binary.cpp69
-rw-r--r--lib/Object/COFFImportFile.cpp14
-rw-r--r--lib/Object/COFFObjectFile.cpp2
-rw-r--r--lib/Object/Decompressor.cpp2
-rw-r--r--lib/Object/ELF.cpp34
-rw-r--r--lib/Object/ELFObjectFile.cpp6
-rw-r--r--lib/Object/IRObjectFile.cpp33
-rw-r--r--lib/Object/IRSymtab.cpp47
-rw-r--r--lib/Object/LLVMBuild.txt2
-rw-r--r--lib/Object/MachOObjectFile.cpp12
-rw-r--r--lib/Object/ModuleSymbolTable.cpp6
-rw-r--r--lib/Object/Object.cpp2
-rw-r--r--lib/Object/ObjectFile.cpp61
-rw-r--r--lib/Object/SymbolicFile.cpp64
-rw-r--r--lib/Object/WasmObjectFile.cpp4
-rw-r--r--lib/Object/WindowsResource.cpp518
-rw-r--r--lib/ObjectYAML/CodeViewYAMLDebugSections.cpp551
-rw-r--r--lib/ObjectYAML/CodeViewYAMLSymbols.cpp13
-rw-r--r--lib/ObjectYAML/CodeViewYAMLTypes.cpp28
-rw-r--r--lib/ObjectYAML/DWARFEmitter.cpp2
-rw-r--r--lib/ObjectYAML/DWARFVisitor.h2
-rw-r--r--lib/ObjectYAML/ELFYAML.cpp26
-rw-r--r--lib/ObjectYAML/MachOYAML.cpp4
-rw-r--r--lib/ObjectYAML/ObjectYAML.cpp2
-rw-r--r--lib/ObjectYAML/WasmYAML.cpp2
-rw-r--r--lib/Option/Arg.cpp2
-rw-r--r--lib/Passes/PassBuilder.cpp10
-rw-r--r--lib/ProfileData/Coverage/CoverageMapping.cpp2
-rw-r--r--lib/ProfileData/Coverage/CoverageMappingWriter.cpp2
-rw-r--r--lib/ProfileData/InstrProf.cpp4
-rw-r--r--lib/ProfileData/InstrProfReader.cpp2
-rw-r--r--lib/ProfileData/InstrProfWriter.cpp2
-rw-r--r--lib/ProfileData/SampleProfWriter.cpp2
-rw-r--r--lib/Support/AMDGPUCodeObjectMetadata.cpp218
-rw-r--r--lib/Support/ARMAttributeParser.cpp2
-rw-r--r--lib/Support/ARMBuildAttrs.cpp2
-rw-r--r--lib/Support/Atomic.cpp2
-rw-r--r--lib/Support/CMakeLists.txt2
-rw-r--r--lib/Support/CommandLine.cpp18
-rw-r--r--lib/Support/ConvertUTF.cpp2
-rw-r--r--lib/Support/ConvertUTFWrapper.cpp2
-rw-r--r--lib/Support/Errno.cpp2
-rw-r--r--lib/Support/Error.cpp1
-rw-r--r--lib/Support/FormattedStream.cpp2
-rw-r--r--lib/Support/LockFileManager.cpp6
-rw-r--r--lib/Support/MD5.cpp2
-rw-r--r--lib/Support/Mutex.cpp2
-rw-r--r--lib/Support/Path.cpp176
-rw-r--r--lib/Support/PrettyStackTrace.cpp2
-rw-r--r--lib/Support/Process.cpp2
-rw-r--r--lib/Support/RWMutex.cpp2
-rw-r--r--lib/Support/SHA1.cpp2
-rw-r--r--lib/Support/Signals.cpp18
-rw-r--r--lib/Support/SourceMgr.cpp6
-rw-r--r--lib/Support/SpecialCaseList.cpp2
-rw-r--r--lib/Support/Statistic.cpp2
-rw-r--r--lib/Support/StringExtras.cpp2
-rw-r--r--lib/Support/TargetRegistry.cpp3
-rw-r--r--lib/Support/ThreadLocal.cpp2
-rw-r--r--lib/Support/Timer.cpp2
-rw-r--r--lib/Support/TrigramIndex.cpp2
-rw-r--r--lib/Support/Triple.cpp21
-rw-r--r--lib/Support/Unix/DynamicLibrary.inc7
-rw-r--r--lib/Support/Unix/Path.inc2
-rw-r--r--lib/Support/Unix/Signals.inc2
-rw-r--r--lib/Support/Unix/Threading.inc8
-rw-r--r--lib/Support/Windows/DynamicLibrary.inc2
-rw-r--r--lib/Support/Windows/WindowsSupport.h4
-rw-r--r--lib/Support/YAMLParser.cpp2
-rw-r--r--lib/TableGen/StringMatcher.cpp2
-rw-r--r--lib/Target/AArch64/AArch64AsmPrinter.cpp6
-rw-r--r--lib/Target/AArch64/AArch64DeadRegisterDefinitionsPass.cpp2
-rw-r--r--lib/Target/AArch64/AArch64ExpandPseudoInsts.cpp2
-rw-r--r--lib/Target/AArch64/AArch64FastISel.cpp6
-rw-r--r--lib/Target/AArch64/AArch64ISelDAGToDAG.cpp2
-rw-r--r--lib/Target/AArch64/AArch64ISelLowering.cpp6
-rw-r--r--lib/Target/AArch64/AArch64InstrInfo.cpp2
-rw-r--r--lib/Target/AArch64/AArch64LegalizerInfo.cpp4
-rw-r--r--lib/Target/AArch64/AArch64LoadStoreOptimizer.cpp2
-rw-r--r--lib/Target/AArch64/AArch64PBQPRegAlloc.cpp2
-rw-r--r--lib/Target/AArch64/AArch64RegisterBankInfo.cpp4
-rw-r--r--lib/Target/AArch64/AArch64TargetObjectFile.cpp2
-rw-r--r--lib/Target/AArch64/AArch64TargetTransformInfo.cpp2
-rw-r--r--lib/Target/AArch64/AsmParser/AArch64AsmParser.cpp2
-rw-r--r--lib/Target/AArch64/MCTargetDesc/AArch64AsmBackend.cpp4
-rw-r--r--lib/Target/AArch64/MCTargetDesc/AArch64ELFObjectWriter.cpp2
-rw-r--r--lib/Target/AArch64/MCTargetDesc/AArch64ELFStreamer.cpp2
-rw-r--r--lib/Target/AArch64/MCTargetDesc/AArch64MachObjectWriter.cpp2
-rw-r--r--lib/Target/AMDGPU/AMDGPU.td33
-rw-r--r--lib/Target/AMDGPU/AMDGPUAliasAnalysis.cpp6
-rw-r--r--lib/Target/AMDGPU/AMDGPUAnnotateUniformValues.cpp2
-rw-r--r--lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp12
-rw-r--r--lib/Target/AMDGPU/AMDGPUAsmPrinter.h2
-rw-r--r--lib/Target/AMDGPU/AMDGPUCallLowering.cpp2
-rw-r--r--lib/Target/AMDGPU/AMDGPUCodeGenPrepare.cpp4
-rw-r--r--lib/Target/AMDGPU/AMDGPUISelDAGToDAG.cpp6
-rw-r--r--lib/Target/AMDGPU/AMDGPUISelLowering.cpp2
-rw-r--r--lib/Target/AMDGPU/AMDGPUInstrInfo.h2
-rw-r--r--lib/Target/AMDGPU/AMDGPUInstructionSelector.h2
-rw-r--r--lib/Target/AMDGPU/AMDGPULegalizerInfo.cpp10
-rw-r--r--lib/Target/AMDGPU/AMDGPUMCInstLower.cpp1
-rw-r--r--lib/Target/AMDGPU/AMDGPUMachineCFGStructurizer.cpp2
-rw-r--r--lib/Target/AMDGPU/AMDGPUMachineFunction.h2
-rw-r--r--lib/Target/AMDGPU/AMDGPUPromoteAlloca.cpp20
-rw-r--r--lib/Target/AMDGPU/AMDGPURegAsmNames.inc.cpp353
-rw-r--r--lib/Target/AMDGPU/AMDGPURegisterBankInfo.h1
-rw-r--r--lib/Target/AMDGPU/AMDGPURegisterInfo.cpp1
-rw-r--r--lib/Target/AMDGPU/AMDGPUSubtarget.h17
-rw-r--r--lib/Target/AMDGPU/AMDGPUTargetMachine.cpp18
-rw-r--r--lib/Target/AMDGPU/AMDGPUTargetObjectFile.cpp6
-rw-r--r--lib/Target/AMDGPU/AMDGPUTargetTransformInfo.cpp2
-rw-r--r--lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp106
-rw-r--r--lib/Target/AMDGPU/CMakeLists.txt1
-rw-r--r--lib/Target/AMDGPU/Disassembler/AMDGPUDisassembler.cpp7
-rw-r--r--lib/Target/AMDGPU/Disassembler/AMDGPUDisassembler.h2
-rw-r--r--lib/Target/AMDGPU/GCNHazardRecognizer.cpp2
-rw-r--r--lib/Target/AMDGPU/InstPrinter/AMDGPUInstPrinter.cpp2
-rw-r--r--lib/Target/AMDGPU/MCTargetDesc/AMDGPUAsmBackend.cpp2
-rw-r--r--lib/Target/AMDGPU/MCTargetDesc/AMDGPUCodeObjectMetadataStreamer.cpp197
-rw-r--r--lib/Target/AMDGPU/MCTargetDesc/AMDGPUCodeObjectMetadataStreamer.h2
-rw-r--r--lib/Target/AMDGPU/MCTargetDesc/AMDGPUELFObjectWriter.cpp2
-rw-r--r--lib/Target/AMDGPU/MCTargetDesc/AMDGPUTargetStreamer.cpp32
-rw-r--r--lib/Target/AMDGPU/MCTargetDesc/AMDGPUTargetStreamer.h12
-rw-r--r--lib/Target/AMDGPU/MCTargetDesc/R600MCCodeEmitter.cpp2
-rw-r--r--lib/Target/AMDGPU/Processors.td75
-rw-r--r--lib/Target/AMDGPU/R600ControlFlowFinalizer.cpp4
-rw-r--r--lib/Target/AMDGPU/R600EmitClauseMarkers.cpp2
-rw-r--r--lib/Target/AMDGPU/R600ExpandSpecialInstrs.cpp2
-rw-r--r--lib/Target/AMDGPU/R600FrameLowering.cpp2
-rw-r--r--lib/Target/AMDGPU/R600ISelLowering.cpp2
-rw-r--r--lib/Target/AMDGPU/R600InstrInfo.cpp4
-rw-r--r--lib/Target/AMDGPU/R600MachineScheduler.cpp4
-rw-r--r--lib/Target/AMDGPU/R600Packetizer.cpp2
-rw-r--r--lib/Target/AMDGPU/SIDebuggerInsertNops.cpp2
-rw-r--r--lib/Target/AMDGPU/SIFixSGPRCopies.cpp2
-rw-r--r--lib/Target/AMDGPU/SIFoldOperands.cpp3
-rw-r--r--lib/Target/AMDGPU/SIFrameLowering.cpp2
-rw-r--r--lib/Target/AMDGPU/SIISelLowering.cpp6
-rw-r--r--lib/Target/AMDGPU/SIInsertWaitcnts.cpp2
-rw-r--r--lib/Target/AMDGPU/SIInstrInfo.cpp2
-rw-r--r--lib/Target/AMDGPU/SILowerControlFlow.cpp2
-rw-r--r--lib/Target/AMDGPU/SILowerI1Copies.cpp2
-rw-r--r--lib/Target/AMDGPU/SIMachineFunctionInfo.h2
-rw-r--r--lib/Target/AMDGPU/SIMachineScheduler.cpp2
-rw-r--r--lib/Target/AMDGPU/SIPeepholeSDWA.cpp97
-rw-r--r--lib/Target/AMDGPU/SIRegisterInfo.cpp62
-rw-r--r--lib/Target/AMDGPU/SIRegisterInfo.h4
-rw-r--r--lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp43
-rw-r--r--lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.h7
-rw-r--r--lib/Target/AMDGPU/VOP3Instructions.td9
-rw-r--r--lib/Target/ARM/ARMAsmPrinter.cpp4
-rw-r--r--lib/Target/ARM/ARMBaseInstrInfo.cpp2
-rw-r--r--lib/Target/ARM/ARMBaseRegisterInfo.cpp4
-rw-r--r--lib/Target/ARM/ARMCallLowering.cpp16
-rw-r--r--lib/Target/ARM/ARMConstantIslandPass.cpp2
-rw-r--r--lib/Target/ARM/ARMFastISel.cpp2
-rw-r--r--lib/Target/ARM/ARMFrameLowering.cpp2
-rw-r--r--lib/Target/ARM/ARMISelLowering.cpp9
-rw-r--r--lib/Target/ARM/ARMInstrVFP.td9
-rw-r--r--lib/Target/ARM/ARMInstructionSelector.cpp70
-rw-r--r--lib/Target/ARM/ARMLegalizerInfo.cpp2
-rw-r--r--lib/Target/ARM/ARMLoadStoreOptimizer.cpp2
-rw-r--r--lib/Target/ARM/ARMMCInstLower.cpp2
-rw-r--r--lib/Target/ARM/ARMRegisterBankInfo.cpp33
-rw-r--r--lib/Target/ARM/ARMSubtarget.cpp4
-rw-r--r--lib/Target/ARM/ARMTargetMachine.cpp15
-rw-r--r--lib/Target/ARM/ARMTargetObjectFile.cpp6
-rw-r--r--lib/Target/ARM/AsmParser/ARMAsmParser.cpp19
-rw-r--r--lib/Target/ARM/Disassembler/ARMDisassembler.cpp2
-rw-r--r--lib/Target/ARM/MCTargetDesc/ARMAsmBackend.cpp66
-rw-r--r--lib/Target/ARM/MCTargetDesc/ARMAsmBackendDarwin.h2
-rw-r--r--lib/Target/ARM/MCTargetDesc/ARMELFObjectWriter.cpp2
-rw-r--r--lib/Target/ARM/MCTargetDesc/ARMELFStreamer.cpp4
-rw-r--r--lib/Target/ARM/MCTargetDesc/ARMFixupKinds.h3
-rw-r--r--lib/Target/ARM/MCTargetDesc/ARMMCCodeEmitter.cpp12
-rw-r--r--lib/Target/ARM/MCTargetDesc/ARMMCTargetDesc.cpp2
-rw-r--r--lib/Target/ARM/MCTargetDesc/ARMMachORelocationInfo.cpp2
-rw-r--r--lib/Target/ARM/MCTargetDesc/ARMMachObjectWriter.cpp4
-rw-r--r--lib/Target/ARM/MCTargetDesc/ARMWinCOFFObjectWriter.cpp2
-rw-r--r--lib/Target/ARM/Thumb1FrameLowering.cpp6
-rw-r--r--lib/Target/ARM/Thumb1InstrInfo.cpp2
-rw-r--r--lib/Target/ARM/Thumb2InstrInfo.cpp20
-rw-r--r--lib/Target/ARM/Thumb2SizeReduction.cpp2
-rw-r--r--lib/Target/AVR/AVR.h2
-rw-r--r--lib/Target/AVR/AVRAsmPrinter.cpp2
-rw-r--r--lib/Target/AVR/AVRRegisterInfo.cpp2
-rw-r--r--lib/Target/AVR/AVRSubtarget.cpp2
-rw-r--r--lib/Target/AVR/AVRSubtarget.h3
-rw-r--r--lib/Target/AVR/AVRTargetMachine.cpp4
-rw-r--r--lib/Target/AVR/AVRTargetObjectFile.cpp2
-rw-r--r--lib/Target/AVR/AsmParser/AVRAsmParser.cpp6
-rw-r--r--lib/Target/AVR/Disassembler/AVRDisassembler.cpp4
-rw-r--r--lib/Target/AVR/MCTargetDesc/AVRELFStreamer.cpp2
-rw-r--r--lib/Target/AVR/MCTargetDesc/AVRMCExpr.cpp2
-rw-r--r--lib/Target/AVR/MCTargetDesc/AVRMCTargetDesc.cpp2
-rw-r--r--lib/Target/BPF/BPFAsmPrinter.cpp4
-rw-r--r--lib/Target/BPF/BPFInstrInfo.cpp2
-rw-r--r--lib/Target/BPF/BPFRegisterInfo.cpp6
-rw-r--r--lib/Target/BPF/BPFTargetMachine.cpp6
-rw-r--r--lib/Target/BPF/Disassembler/BPFDisassembler.cpp4
-rw-r--r--lib/Target/BPF/InstPrinter/BPFInstPrinter.cpp2
-rw-r--r--lib/Target/BPF/MCTargetDesc/BPFELFObjectWriter.cpp2
-rw-r--r--lib/Target/BPF/MCTargetDesc/BPFMCTargetDesc.cpp2
-rw-r--r--lib/Target/BPF/MCTargetDesc/BPFMCTargetDesc.h2
-rw-r--r--lib/Target/Hexagon/AsmParser/HexagonAsmParser.cpp6
-rw-r--r--lib/Target/Hexagon/BitTracker.cpp2
-rw-r--r--lib/Target/Hexagon/Disassembler/HexagonDisassembler.cpp6
-rw-r--r--lib/Target/Hexagon/HexagonAsmPrinter.cpp4
-rw-r--r--lib/Target/Hexagon/HexagonBitSimplify.cpp2
-rw-r--r--lib/Target/Hexagon/HexagonBitTracker.cpp2
-rw-r--r--lib/Target/Hexagon/HexagonBlockRanges.h2
-rw-r--r--lib/Target/Hexagon/HexagonCommonGEP.cpp16
-rw-r--r--lib/Target/Hexagon/HexagonConstPropagation.cpp2
-rw-r--r--lib/Target/Hexagon/HexagonCopyToCombine.cpp2
-rw-r--r--lib/Target/Hexagon/HexagonEarlyIfConv.cpp2
-rw-r--r--lib/Target/Hexagon/HexagonFixupHwLoops.cpp3
-rw-r--r--lib/Target/Hexagon/HexagonFrameLowering.cpp2
-rw-r--r--lib/Target/Hexagon/HexagonGenExtract.cpp2
-rw-r--r--lib/Target/Hexagon/HexagonGenInsert.cpp4
-rw-r--r--lib/Target/Hexagon/HexagonGenMux.cpp7
-rw-r--r--lib/Target/Hexagon/HexagonISelLowering.cpp4
-rw-r--r--lib/Target/Hexagon/HexagonInstrInfo.cpp6
-rw-r--r--lib/Target/Hexagon/HexagonLoopIdiomRecognition.cpp4
-rw-r--r--lib/Target/Hexagon/HexagonMachineScheduler.cpp57
-rw-r--r--lib/Target/Hexagon/HexagonPatterns.td89
-rw-r--r--lib/Target/Hexagon/HexagonSplitDouble.cpp2
-rw-r--r--lib/Target/Hexagon/HexagonTargetMachine.cpp36
-rw-r--r--lib/Target/Hexagon/HexagonTargetObjectFile.cpp2
-rw-r--r--lib/Target/Hexagon/HexagonVLIWPacketizer.cpp2
-rw-r--r--lib/Target/Hexagon/MCTargetDesc/HexagonAsmBackend.cpp2
-rw-r--r--lib/Target/Hexagon/MCTargetDesc/HexagonInstPrinter.cpp2
-rw-r--r--lib/Target/Hexagon/MCTargetDesc/HexagonMCCodeEmitter.cpp2
-rw-r--r--lib/Target/Hexagon/MCTargetDesc/HexagonMCELFStreamer.cpp2
-rw-r--r--lib/Target/Hexagon/MCTargetDesc/HexagonMCShuffler.cpp2
-rw-r--r--lib/Target/Hexagon/MCTargetDesc/HexagonMCTargetDesc.cpp6
-rw-r--r--lib/Target/Hexagon/RDFDeadCode.cpp2
-rw-r--r--lib/Target/Hexagon/RDFGraph.cpp2
-rw-r--r--lib/Target/Hexagon/RDFLiveness.cpp2
-rw-r--r--lib/Target/Lanai/AsmParser/LanaiAsmParser.cpp2
-rw-r--r--lib/Target/Lanai/LanaiTargetObjectFile.cpp2
-rw-r--r--lib/Target/Lanai/MCTargetDesc/LanaiELFObjectWriter.cpp2
-rw-r--r--lib/Target/Lanai/MCTargetDesc/LanaiMCCodeEmitter.cpp2
-rw-r--r--lib/Target/Lanai/MCTargetDesc/LanaiMCTargetDesc.cpp2
-rw-r--r--lib/Target/MSP430/MSP430AsmPrinter.cpp2
-rw-r--r--lib/Target/Mips/AsmParser/MipsAsmParser.cpp10
-rw-r--r--lib/Target/Mips/Disassembler/MipsDisassembler.cpp4
-rw-r--r--lib/Target/Mips/MCTargetDesc/MipsAsmBackend.cpp2
-rw-r--r--lib/Target/Mips/MCTargetDesc/MipsELFObjectWriter.cpp2
-rw-r--r--lib/Target/Mips/MCTargetDesc/MipsELFStreamer.cpp2
-rw-r--r--lib/Target/Mips/MCTargetDesc/MipsMCCodeEmitter.cpp2
-rw-r--r--lib/Target/Mips/MCTargetDesc/MipsMCExpr.cpp4
-rw-r--r--lib/Target/Mips/MCTargetDesc/MipsOptionRecord.cpp4
-rw-r--r--lib/Target/Mips/MCTargetDesc/MipsTargetStreamer.cpp4
-rw-r--r--lib/Target/Mips/Mips.td3
-rw-r--r--lib/Target/Mips/Mips16FrameLowering.cpp4
-rw-r--r--lib/Target/Mips/MipsAsmPrinter.cpp4
-rw-r--r--lib/Target/Mips/MipsCCState.cpp61
-rw-r--r--lib/Target/Mips/MipsCCState.h33
-rw-r--r--lib/Target/Mips/MipsCallingConv.td10
-rw-r--r--lib/Target/Mips/MipsConstantIslandPass.cpp2
-rw-r--r--lib/Target/Mips/MipsDSPInstrInfo.td5
-rw-r--r--lib/Target/Mips/MipsFastISel.cpp2
-rw-r--r--lib/Target/Mips/MipsISelLowering.cpp81
-rw-r--r--lib/Target/Mips/MipsISelLowering.h27
-rw-r--r--lib/Target/Mips/MipsInstrFPU.td18
-rw-r--r--lib/Target/Mips/MipsInstrInfo.td6
-rw-r--r--lib/Target/Mips/MipsMachineFunction.cpp8
-rw-r--r--lib/Target/Mips/MipsOptimizePICCall.cpp2
-rw-r--r--lib/Target/Mips/MipsOs16.cpp2
-rw-r--r--lib/Target/Mips/MipsRegisterInfo.cpp4
-rw-r--r--lib/Target/Mips/MipsSEFrameLowering.cpp2
-rw-r--r--lib/Target/Mips/MipsSEISelDAGToDAG.cpp2
-rw-r--r--lib/Target/Mips/MipsSubtarget.cpp6
-rw-r--r--lib/Target/Mips/MipsSubtarget.h5
-rw-r--r--lib/Target/Mips/MipsTargetMachine.cpp2
-rw-r--r--lib/Target/Mips/MipsTargetObjectFile.cpp2
-rw-r--r--lib/Target/NVPTX/NVPTXAsmPrinter.cpp4
-rw-r--r--lib/Target/NVPTX/NVPTXGenericToNVVM.cpp2
-rw-r--r--lib/Target/NVPTX/NVPTXISelLowering.cpp2
-rw-r--r--lib/Target/NVPTX/NVPTXInstrInfo.cpp2
-rw-r--r--lib/Target/NVPTX/NVPTXLowerArgs.cpp2
-rw-r--r--lib/Target/NVPTX/NVPTXPeephole.cpp2
-rw-r--r--lib/Target/NVPTX/NVPTXTargetMachine.cpp2
-rw-r--r--lib/Target/NVPTX/NVVMIntrRange.cpp2
-rw-r--r--lib/Target/PowerPC/InstPrinter/PPCInstPrinter.cpp2
-rw-r--r--lib/Target/PowerPC/MCTargetDesc/PPCAsmBackend.cpp6
-rw-r--r--lib/Target/PowerPC/MCTargetDesc/PPCELFObjectWriter.cpp2
-rw-r--r--lib/Target/PowerPC/MCTargetDesc/PPCMCExpr.cpp2
-rw-r--r--lib/Target/PowerPC/MCTargetDesc/PPCMCTargetDesc.cpp6
-rw-r--r--lib/Target/PowerPC/MCTargetDesc/PPCMachObjectWriter.cpp8
-rw-r--r--lib/Target/PowerPC/PPCAsmPrinter.cpp10
-rw-r--r--lib/Target/PowerPC/PPCBoolRetToInt.cpp42
-rw-r--r--lib/Target/PowerPC/PPCBranchSelector.cpp2
-rw-r--r--lib/Target/PowerPC/PPCCTRLoops.cpp2
-rw-r--r--lib/Target/PowerPC/PPCEarlyReturn.cpp2
-rw-r--r--lib/Target/PowerPC/PPCFastISel.cpp4
-rw-r--r--lib/Target/PowerPC/PPCISelDAGToDAG.cpp65
-rw-r--r--lib/Target/PowerPC/PPCISelLowering.cpp10
-rw-r--r--lib/Target/PowerPC/PPCInstrVSX.td51
-rw-r--r--lib/Target/PowerPC/PPCMCInstLower.cpp2
-rw-r--r--lib/Target/PowerPC/PPCMIPeephole.cpp2
-rw-r--r--lib/Target/PowerPC/PPCTLSDynamicCall.cpp2
-rw-r--r--lib/Target/PowerPC/PPCTOCRegDeps.cpp2
-rw-r--r--lib/Target/PowerPC/PPCTargetMachine.cpp2
-rw-r--r--lib/Target/PowerPC/PPCVSXCopy.cpp2
-rw-r--r--lib/Target/PowerPC/PPCVSXFMAMutate.cpp2
-rw-r--r--lib/Target/PowerPC/PPCVSXSwapRemoval.cpp2
-rw-r--r--lib/Target/RISCV/MCTargetDesc/RISCVAsmBackend.cpp2
-rw-r--r--lib/Target/RISCV/MCTargetDesc/RISCVMCCodeEmitter.cpp2
-rw-r--r--lib/Target/RISCV/MCTargetDesc/RISCVMCTargetDesc.h2
-rw-r--r--lib/Target/RISCV/RISCVTargetMachine.cpp2
-rw-r--r--lib/Target/Sparc/AsmParser/SparcAsmParser.cpp4
-rw-r--r--lib/Target/Sparc/Disassembler/SparcDisassembler.cpp4
-rw-r--r--lib/Target/Sparc/MCTargetDesc/SparcAsmBackend.cpp2
-rw-r--r--lib/Target/Sparc/MCTargetDesc/SparcMCAsmInfo.cpp2
-rw-r--r--lib/Target/Sparc/MCTargetDesc/SparcMCExpr.cpp1
-rw-r--r--lib/Target/Sparc/SparcAsmPrinter.cpp2
-rw-r--r--lib/Target/Sparc/SparcMCInstLower.cpp2
-rw-r--r--lib/Target/Sparc/SparcTargetMachine.cpp4
-rw-r--r--lib/Target/Sparc/SparcTargetObjectFile.cpp2
-rw-r--r--lib/Target/SystemZ/AsmParser/SystemZAsmParser.cpp2
-rw-r--r--lib/Target/SystemZ/MCTargetDesc/SystemZMCAsmBackend.cpp2
-rw-r--r--lib/Target/SystemZ/MCTargetDesc/SystemZMCObjectWriter.cpp2
-rw-r--r--lib/Target/SystemZ/SystemZHazardRecognizer.h2
-rw-r--r--lib/Target/SystemZ/SystemZISelLowering.cpp25
-rw-r--r--lib/Target/SystemZ/SystemZInstrInfo.cpp2
-rw-r--r--lib/Target/SystemZ/SystemZLDCleanup.cpp2
-rw-r--r--lib/Target/SystemZ/SystemZRegisterInfo.cpp2
-rw-r--r--lib/Target/SystemZ/SystemZShortenInst.cpp2
-rw-r--r--lib/Target/SystemZ/SystemZSubtarget.h2
-rw-r--r--lib/Target/SystemZ/SystemZTDC.cpp4
-rw-r--r--lib/Target/SystemZ/SystemZTargetMachine.cpp4
-rw-r--r--lib/Target/Target.cpp4
-rw-r--r--lib/Target/TargetLoweringObjectFile.cpp16
-rw-r--r--lib/Target/TargetMachineC.cpp4
-rw-r--r--lib/Target/WebAssembly/Disassembler/WebAssemblyDisassembler.cpp2
-rw-r--r--lib/Target/WebAssembly/InstPrinter/WebAssemblyInstPrinter.h2
-rw-r--r--lib/Target/WebAssembly/MCTargetDesc/WebAssemblyAsmBackend.cpp2
-rw-r--r--lib/Target/WebAssembly/MCTargetDesc/WebAssemblyMCCodeEmitter.cpp2
-rw-r--r--lib/Target/WebAssembly/MCTargetDesc/WebAssemblyMCTargetDesc.h2
-rw-r--r--lib/Target/WebAssembly/MCTargetDesc/WebAssemblyTargetStreamer.h2
-rw-r--r--lib/Target/WebAssembly/MCTargetDesc/WebAssemblyWasmObjectWriter.cpp11
-rw-r--r--lib/Target/WebAssembly/WebAssemblyCFGSort.cpp2
-rw-r--r--lib/Target/WebAssembly/WebAssemblyCFGStackify.cpp2
-rw-r--r--lib/Target/WebAssembly/WebAssemblyCallIndirectFixup.cpp2
-rw-r--r--lib/Target/WebAssembly/WebAssemblyFastISel.cpp2
-rw-r--r--lib/Target/WebAssembly/WebAssemblyFixIrreducibleControlFlow.cpp2
-rw-r--r--lib/Target/WebAssembly/WebAssemblyISelDAGToDAG.cpp2
-rw-r--r--lib/Target/WebAssembly/WebAssemblyLowerBrUnless.cpp2
-rw-r--r--lib/Target/WebAssembly/WebAssemblyPrepareForLiveIntervals.cpp2
-rw-r--r--lib/Target/WebAssembly/WebAssemblyRegNumbering.cpp2
-rw-r--r--lib/Target/WebAssembly/WebAssemblyRegStackify.cpp2
-rw-r--r--lib/Target/WebAssembly/WebAssemblyReplacePhysRegs.cpp2
-rw-r--r--lib/Target/WebAssembly/WebAssemblySetP2AlignOperands.cpp2
-rw-r--r--lib/Target/WebAssembly/WebAssemblyStoreResults.cpp2
-rw-r--r--lib/Target/WebAssembly/WebAssemblyTargetMachine.cpp4
-rw-r--r--lib/Target/X86/AsmParser/X86AsmInstrumentation.cpp4
-rw-r--r--lib/Target/X86/AsmParser/X86Operand.h2
-rw-r--r--lib/Target/X86/CMakeLists.txt1
-rw-r--r--lib/Target/X86/Disassembler/X86Disassembler.cpp2
-rw-r--r--lib/Target/X86/Disassembler/X86DisassemblerDecoder.cpp8
-rw-r--r--lib/Target/X86/InstPrinter/X86ATTInstPrinter.cpp2
-rw-r--r--lib/Target/X86/InstPrinter/X86InstComments.cpp2
-rw-r--r--lib/Target/X86/InstPrinter/X86IntelInstPrinter.cpp2
-rw-r--r--lib/Target/X86/MCTargetDesc/X86AsmBackend.cpp4
-rw-r--r--lib/Target/X86/MCTargetDesc/X86ELFObjectWriter.cpp2
-rw-r--r--lib/Target/X86/MCTargetDesc/X86MCAsmInfo.cpp2
-rw-r--r--lib/Target/X86/MCTargetDesc/X86MachObjectWriter.cpp4
-rw-r--r--lib/Target/X86/MCTargetDesc/X86WinCOFFObjectWriter.cpp2
-rw-r--r--lib/Target/X86/X86AsmPrinter.cpp2
-rw-r--r--lib/Target/X86/X86FastISel.cpp6
-rw-r--r--lib/Target/X86/X86FrameLowering.cpp2
-rw-r--r--lib/Target/X86/X86ISelLowering.cpp193
-rw-r--r--lib/Target/X86/X86InstrFragmentsSIMD.td41
-rw-r--r--lib/Target/X86/X86InstrInfo.cpp3400
-rw-r--r--lib/Target/X86/X86InstrSSE.td10
-rw-r--r--lib/Target/X86/X86MCInstLower.cpp16
-rw-r--r--lib/Target/X86/X86OptimizeLEAs.cpp2
-rw-r--r--lib/Target/X86/X86SchedHaswell.td33
-rw-r--r--lib/Target/X86/X86SchedSandyBridge.td25
-rw-r--r--lib/Target/X86/X86Schedule.td4
-rw-r--r--lib/Target/X86/X86ScheduleBtVer2.td32
-rw-r--r--lib/Target/X86/X86ScheduleSLM.td27
-rw-r--r--lib/Target/X86/X86SelectionDAGInfo.cpp4
-rw-r--r--lib/Target/X86/X86Subtarget.cpp2
-rw-r--r--lib/Target/X86/X86TargetMachine.cpp2
-rw-r--r--lib/Target/X86/X86TargetObjectFile.cpp4
-rw-r--r--lib/Target/X86/X86TargetTransformInfo.cpp13
-rw-r--r--lib/Target/X86/X86TargetTransformInfo.h4
-rw-r--r--lib/Target/X86/X86WinEHState.cpp2
-rw-r--r--lib/Target/XCore/MCTargetDesc/XCoreMCTargetDesc.cpp4
-rw-r--r--lib/Target/XCore/XCoreAsmPrinter.cpp2
-rw-r--r--lib/Target/XCore/XCoreTargetMachine.cpp2
-rw-r--r--lib/Target/XCore/XCoreTargetMachine.h2
-rw-r--r--lib/Target/XCore/XCoreTargetObjectFile.cpp2
-rw-r--r--lib/ToolDrivers/llvm-lib/LLVMBuild.txt2
-rw-r--r--lib/ToolDrivers/llvm-lib/LibDriver.cpp12
-rw-r--r--lib/Transforms/Coroutines/CoroSplit.cpp2
-rw-r--r--lib/Transforms/IPO/ElimAvailExtern.cpp2
-rw-r--r--lib/Transforms/IPO/ExtractGV.cpp14
-rw-r--r--lib/Transforms/IPO/FunctionAttrs.cpp3
-rw-r--r--lib/Transforms/IPO/GlobalSplit.cpp2
-rw-r--r--lib/Transforms/IPO/IPConstantPropagation.cpp2
-rw-r--r--lib/Transforms/IPO/IPO.cpp4
-rw-r--r--lib/Transforms/IPO/InferFunctionAttrs.cpp2
-rw-r--r--lib/Transforms/IPO/Inliner.cpp8
-rw-r--r--lib/Transforms/IPO/LoopExtractor.cpp2
-rw-r--r--lib/Transforms/IPO/LowerTypeTests.cpp56
-rw-r--r--lib/Transforms/IPO/PruneEH.cpp4
-rw-r--r--lib/Transforms/IPO/SampleProfile.cpp7
-rw-r--r--lib/Transforms/IPO/StripSymbols.cpp2
-rw-r--r--lib/Transforms/IPO/ThinLTOBitcodeWriter.cpp12
-rw-r--r--lib/Transforms/IPO/WholeProgramDevirt.cpp2
-rw-r--r--lib/Transforms/InstCombine/InstCombineAddSub.cpp23
-rw-r--r--lib/Transforms/InstCombine/InstCombineAndOrXor.cpp27
-rw-r--r--lib/Transforms/InstCombine/InstCombineCalls.cpp16
-rw-r--r--lib/Transforms/InstCombine/InstCombineCasts.cpp8
-rw-r--r--lib/Transforms/InstCombine/InstCombineCompares.cpp70
-rw-r--r--lib/Transforms/InstCombine/InstCombineInternal.h2
-rw-r--r--lib/Transforms/InstCombine/InstCombineMulDivRem.cpp21
-rw-r--r--lib/Transforms/InstCombine/InstCombinePHI.cpp4
-rw-r--r--lib/Transforms/InstCombine/InstCombineSelect.cpp7
-rw-r--r--lib/Transforms/InstCombine/InstCombineShifts.cpp30
-rw-r--r--lib/Transforms/InstCombine/InstCombineSimplifyDemanded.cpp8
-rw-r--r--lib/Transforms/InstCombine/InstCombineVectorOps.cpp9
-rw-r--r--lib/Transforms/InstCombine/InstructionCombining.cpp45
-rw-r--r--lib/Transforms/Instrumentation/BoundsChecking.cpp2
-rw-r--r--lib/Transforms/Instrumentation/DataFlowSanitizer.cpp4
-rw-r--r--lib/Transforms/Instrumentation/EfficiencySanitizer.cpp2
-rw-r--r--lib/Transforms/Instrumentation/InstrProfiling.cpp2
-rw-r--r--lib/Transforms/Instrumentation/ThreadSanitizer.cpp2
-rw-r--r--lib/Transforms/ObjCARC/BlotMapVector.h2
-rw-r--r--lib/Transforms/ObjCARC/DependencyAnalysis.cpp2
-rw-r--r--lib/Transforms/ObjCARC/ObjCARCContract.cpp2
-rw-r--r--lib/Transforms/ObjCARC/ObjCARCOpts.cpp2
-rw-r--r--lib/Transforms/ObjCARC/ProvenanceAnalysis.cpp2
-rw-r--r--lib/Transforms/ObjCARC/ProvenanceAnalysisEvaluator.cpp4
-rw-r--r--lib/Transforms/ObjCARC/PtrState.h2
-rw-r--r--lib/Transforms/Scalar/AlignmentFromAssumptions.cpp4
-rw-r--r--lib/Transforms/Scalar/ConstantProp.cpp6
-rw-r--r--lib/Transforms/Scalar/DCE.cpp2
-rw-r--r--lib/Transforms/Scalar/FlattenCFGPass.cpp2
-rw-r--r--lib/Transforms/Scalar/GVNHoist.cpp2
-rw-r--r--lib/Transforms/Scalar/GVNSink.cpp6
-rw-r--r--lib/Transforms/Scalar/GuardWidening.cpp2
-rw-r--r--lib/Transforms/Scalar/IndVarSimplify.cpp7
-rw-r--r--lib/Transforms/Scalar/InductiveRangeCheckElimination.cpp35
-rw-r--r--lib/Transforms/Scalar/InferAddressSpaces.cpp9
-rw-r--r--lib/Transforms/Scalar/JumpThreading.cpp10
-rw-r--r--lib/Transforms/Scalar/LoadCombine.cpp2
-rw-r--r--lib/Transforms/Scalar/LoopIdiomRecognize.cpp77
-rw-r--r--lib/Transforms/Scalar/LoopPredication.cpp2
-rw-r--r--lib/Transforms/Scalar/LoopRerollPass.cpp4
-rw-r--r--lib/Transforms/Scalar/LoopStrengthReduce.cpp167
-rw-r--r--lib/Transforms/Scalar/LoopUnswitch.cpp16
-rw-r--r--lib/Transforms/Scalar/LowerExpectIntrinsic.cpp4
-rw-r--r--lib/Transforms/Scalar/LowerGuardIntrinsic.cpp2
-rw-r--r--lib/Transforms/Scalar/MemCpyOptimizer.cpp8
-rw-r--r--lib/Transforms/Scalar/NewGVN.cpp80
-rw-r--r--lib/Transforms/Scalar/Reg2Mem.cpp2
-rw-r--r--lib/Transforms/Scalar/RewriteStatepointsForGC.cpp18
-rw-r--r--lib/Transforms/Scalar/SCCP.cpp4
-rw-r--r--lib/Transforms/Scalar/SROA.cpp19
-rw-r--r--lib/Transforms/Scalar/Scalar.cpp6
-rw-r--r--lib/Transforms/Scalar/Scalarizer.cpp2
-rw-r--r--lib/Transforms/Scalar/SeparateConstOffsetFromGEP.cpp10
-rw-r--r--lib/Transforms/Scalar/SimpleLoopUnswitch.cpp4
-rw-r--r--lib/Transforms/Scalar/Sink.cpp2
-rw-r--r--lib/Transforms/Scalar/StructurizeCFG.cpp2
-rw-r--r--lib/Transforms/Scalar/TailRecursionElimination.cpp4
-rw-r--r--lib/Transforms/Utils/CMakeLists.txt1
-rw-r--r--lib/Transforms/Utils/CloneFunction.cpp2
-rw-r--r--lib/Transforms/Utils/CloneModule.cpp4
-rw-r--r--lib/Transforms/Utils/DemoteRegToStack.cpp2
-rw-r--r--lib/Transforms/Utils/Evaluator.cpp4
-rw-r--r--lib/Transforms/Utils/FlattenCFG.cpp2
-rw-r--r--lib/Transforms/Utils/FunctionComparator.cpp2
-rw-r--r--lib/Transforms/Utils/FunctionImportUtils.cpp2
-rw-r--r--lib/Transforms/Utils/GlobalStatus.cpp2
-rw-r--r--lib/Transforms/Utils/InlineFunction.cpp6
-rw-r--r--lib/Transforms/Utils/InstructionNamer.cpp2
-rw-r--r--lib/Transforms/Utils/Local.cpp2
-rw-r--r--lib/Transforms/Utils/LoopSimplify.cpp4
-rw-r--r--lib/Transforms/Utils/LoopUnroll.cpp2
-rw-r--r--lib/Transforms/Utils/LoopUnrollRuntime.cpp2
-rw-r--r--lib/Transforms/Utils/LoopUtils.cpp3
-rw-r--r--lib/Transforms/Utils/LowerMemIntrinsics.cpp4
-rw-r--r--lib/Transforms/Utils/LowerSwitch.cpp2
-rw-r--r--lib/Transforms/Utils/MetaRenamer.cpp2
-rw-r--r--lib/Transforms/Utils/OrderedInstructions.cpp33
-rw-r--r--lib/Transforms/Utils/SSAUpdater.cpp4
-rw-r--r--lib/Transforms/Utils/SanitizerStats.cpp2
-rw-r--r--lib/Transforms/Utils/SimplifyCFG.cpp5
-rw-r--r--lib/Transforms/Utils/SimplifyInstructions.cpp2
-rw-r--r--lib/Transforms/Utils/SimplifyLibCalls.cpp2
-rw-r--r--lib/Transforms/Utils/StripGCRelocates.cpp2
-rw-r--r--lib/Transforms/Utils/StripNonLineTableDebugInfo.cpp2
-rw-r--r--lib/Transforms/Utils/SymbolRewriter.cpp2
-rw-r--r--lib/Transforms/Utils/Utils.cpp2
-rw-r--r--lib/Transforms/Vectorize/BBVectorize.cpp2
-rw-r--r--lib/Transforms/Vectorize/LoopVectorize.cpp9
-rw-r--r--lib/Transforms/Vectorize/Vectorize.cpp2
-rw-r--r--lib/XRay/InstrumentationMap.cpp2
-rw-r--r--test/Analysis/BranchProbabilityInfo/libfunc_call.ll264
-rw-r--r--test/Analysis/ConstantFolding/gep-constanfolding-error.ll2
-rw-r--r--test/Analysis/LazyValueAnalysis/lvi-after-jumpthreading.ll68
-rw-r--r--test/Bindings/OCaml/core.ml11
-rw-r--r--test/Bitcode/ptest-old.ll2
-rw-r--r--test/BugPoint/unsymbolized.ll21
-rw-r--r--test/CMakeLists.txt2
-rw-r--r--test/CodeGen/AArch64/GlobalISel/arm64-irtranslator-stackprotect.ll2
-rw-r--r--test/CodeGen/AArch64/GlobalISel/arm64-irtranslator.ll23
-rw-r--r--test/CodeGen/AArch64/GlobalISel/arm64-regbankselect.mir158
-rw-r--r--test/CodeGen/AArch64/GlobalISel/call-translator-ios.ll4
-rw-r--r--test/CodeGen/AArch64/GlobalISel/call-translator.ll10
-rw-r--r--test/CodeGen/AArch64/GlobalISel/debug-insts.ll4
-rw-r--r--test/CodeGen/AArch64/GlobalISel/localizer-in-O0-pipeline.mir16
-rw-r--r--test/CodeGen/AArch64/GlobalISel/localizer.mir80
-rw-r--r--test/CodeGen/AArch64/GlobalISel/regbankselect-dbg-value.mir2
-rw-r--r--test/CodeGen/AArch64/GlobalISel/regbankselect-default.mir164
-rw-r--r--test/CodeGen/AArch64/GlobalISel/select-binop.mir200
-rw-r--r--test/CodeGen/AArch64/GlobalISel/select-bitcast.mir32
-rw-r--r--test/CodeGen/AArch64/GlobalISel/select-fp-casts.mir72
-rw-r--r--test/CodeGen/AArch64/GlobalISel/select-int-ext.mir46
-rw-r--r--test/CodeGen/AArch64/GlobalISel/select-int-ptr-casts.mir24
-rw-r--r--test/CodeGen/AArch64/GlobalISel/select-load.mir100
-rw-r--r--test/CodeGen/AArch64/GlobalISel/select-muladd.mir14
-rw-r--r--test/CodeGen/AArch64/GlobalISel/select-store.mir84
-rw-r--r--test/CodeGen/AArch64/GlobalISel/select-trunc.mir12
-rw-r--r--test/CodeGen/AArch64/GlobalISel/select-xor.mir30
-rw-r--r--test/CodeGen/AArch64/GlobalISel/select.mir52
-rw-r--r--test/CodeGen/AArch64/GlobalISel/varargs-ios-translator.ll2
-rw-r--r--test/CodeGen/AArch64/arm64-fast-isel-conversion-fallback.ll131
-rw-r--r--test/CodeGen/AArch64/spill-undef.mir67
-rw-r--r--test/CodeGen/AMDGPU/GlobalISel/legalize-icmp.mir24
-rw-r--r--test/CodeGen/AMDGPU/GlobalISel/legalize-select.mir28
-rw-r--r--test/CodeGen/AMDGPU/GlobalISel/regbankselect.mir14
-rw-r--r--test/CodeGen/AMDGPU/add.v2i16.ll8
-rw-r--r--test/CodeGen/AMDGPU/ashr.v2i16.ll2
-rw-r--r--test/CodeGen/AMDGPU/branch-relax-spill.ll420
-rw-r--r--test/CodeGen/AMDGPU/clamp-omod-special-case.mir20
-rw-r--r--test/CodeGen/AMDGPU/exceed-max-sgprs.ll142
-rw-r--r--test/CodeGen/AMDGPU/fabs.f16.ll8
-rw-r--r--test/CodeGen/AMDGPU/fadd.f16.ll6
-rw-r--r--test/CodeGen/AMDGPU/fcanonicalize.f16.ll6
-rw-r--r--test/CodeGen/AMDGPU/flat-scratch-reg.ll14
-rw-r--r--test/CodeGen/AMDGPU/fmul.f16.ll6
-rw-r--r--test/CodeGen/AMDGPU/fneg-fabs.f16.ll8
-rw-r--r--test/CodeGen/AMDGPU/fneg.f16.ll2
-rw-r--r--test/CodeGen/AMDGPU/fptosi.f16.ll2
-rw-r--r--test/CodeGen/AMDGPU/fptoui.f16.ll2
-rw-r--r--test/CodeGen/AMDGPU/fsub.f16.ll4
-rw-r--r--test/CodeGen/AMDGPU/hsa-note-no-func.ll13
-rw-r--r--test/CodeGen/AMDGPU/illegal-sgpr-to-vgpr-copy.ll20
-rw-r--r--test/CodeGen/AMDGPU/immv216.ll38
-rw-r--r--test/CodeGen/AMDGPU/indirect-addressing-si.ll2
-rw-r--r--test/CodeGen/AMDGPU/inline-asm.ll12
-rw-r--r--test/CodeGen/AMDGPU/insert_vector_elt.v2i16.ll8
-rw-r--r--test/CodeGen/AMDGPU/limit-coalesce.mir14
-rw-r--r--test/CodeGen/AMDGPU/llvm.amdgcn.alignb.ll23
-rw-r--r--test/CodeGen/AMDGPU/llvm.amdgcn.mqsad.pk.u16.u8.ll22
-rw-r--r--test/CodeGen/AMDGPU/llvm.amdgcn.mqsad.u32.u8.ll57
-rw-r--r--test/CodeGen/AMDGPU/llvm.amdgcn.qsad.pk.u16.u8.ll22
-rw-r--r--test/CodeGen/AMDGPU/llvm.fmuladd.f16.ll2
-rw-r--r--test/CodeGen/AMDGPU/llvm.maxnum.f16.ll6
-rw-r--r--test/CodeGen/AMDGPU/llvm.minnum.f16.ll6
-rw-r--r--test/CodeGen/AMDGPU/partial-sgpr-to-vgpr-spills.ll10
-rw-r--r--test/CodeGen/AMDGPU/promote-alloca-array-aggregate.ll131
-rw-r--r--test/CodeGen/AMDGPU/rename-independent-subregs-invalid-mac-operands.mir69
-rw-r--r--test/CodeGen/AMDGPU/scratch-simple.ll6
-rw-r--r--test/CodeGen/AMDGPU/sdwa-peephole.ll24
-rw-r--r--test/CodeGen/AMDGPU/shl.v2i16.ll2
-rw-r--r--test/CodeGen/AMDGPU/si-spill-sgpr-stack.ll88
-rw-r--r--test/CodeGen/AMDGPU/skip-if-dead.ll12
-rw-r--r--test/CodeGen/AMDGPU/sminmax.v2i16.ll8
-rw-r--r--test/CodeGen/AMDGPU/spill-scavenge-offset.ll14
-rw-r--r--test/CodeGen/AMDGPU/sub.v2i16.ll12
-rw-r--r--test/CodeGen/AMDGPU/undefined-subreg-liverange.ll6
-rw-r--r--test/CodeGen/AMDGPU/v_mac_f16.ll31
-rw-r--r--test/CodeGen/ARM/GlobalISel/arm-instruction-select-combos.mir149
-rw-r--r--test/CodeGen/ARM/GlobalISel/arm-instruction-select.mir115
-rw-r--r--test/CodeGen/ARM/GlobalISel/arm-irtranslator.ll90
-rw-r--r--test/CodeGen/ARM/GlobalISel/arm-isel.ll84
-rw-r--r--test/CodeGen/ARM/GlobalISel/arm-legalize-fp.mir36
-rw-r--r--test/CodeGen/ARM/GlobalISel/arm-legalizer.mir242
-rw-r--r--test/CodeGen/ARM/GlobalISel/arm-regbankselect.mir281
-rw-r--r--test/CodeGen/ARM/clang-section.ll140
-rw-r--r--test/CodeGen/ARM/cortex-a57-misched-vfma.ll91
-rw-r--r--test/CodeGen/ARM/invalidated-save-point.ll4
-rw-r--r--test/CodeGen/Generic/llc-start-stop.ll4
-rw-r--r--test/CodeGen/Hexagon/common-gep-inbounds.ll20
-rw-r--r--test/CodeGen/Hexagon/mux-undef.ll27
-rw-r--r--test/CodeGen/MIR/AArch64/generic-virtual-registers-error.mir1
-rw-r--r--test/CodeGen/MIR/AArch64/generic-virtual-registers-with-regbank-error.mir1
-rw-r--r--test/CodeGen/MIR/AArch64/register-operand-bank.mir4
-rw-r--r--test/CodeGen/MIR/AArch64/stack-object-local-offset.mir4
-rw-r--r--test/CodeGen/MIR/Generic/frame-info.mir5
-rw-r--r--test/CodeGen/MIR/Generic/function-missing-machine-function.mir13
-rw-r--r--test/CodeGen/MIR/X86/callee-saved-info.mir4
-rw-r--r--test/CodeGen/MIR/X86/empty0.mir6
-rw-r--r--test/CodeGen/MIR/X86/empty1.mir8
-rw-r--r--test/CodeGen/MIR/X86/empty2.mir8
-rw-r--r--test/CodeGen/MIR/X86/fixed-stack-objects.mir2
-rw-r--r--test/CodeGen/MIR/X86/generic-instr-type.mir10
-rw-r--r--test/CodeGen/MIR/X86/inline-asm.mir12
-rw-r--r--test/CodeGen/MIR/X86/register-operand-class.mir12
-rw-r--r--test/CodeGen/MIR/X86/roundtrip.mir20
-rw-r--r--test/CodeGen/MIR/X86/simple-register-allocation-hints.mir2
-rw-r--r--test/CodeGen/MIR/X86/spill-slot-fixed-stack-objects.mir2
-rw-r--r--test/CodeGen/MIR/X86/stack-object-debug-info.mir5
-rw-r--r--test/CodeGen/MIR/X86/stack-objects.mir9
-rw-r--r--test/CodeGen/MIR/X86/variable-sized-stack-objects.mir8
-rw-r--r--test/CodeGen/MIR/X86/virtual-registers.mir12
-rw-r--r--test/CodeGen/Mips/biggot.ll3
-rw-r--r--test/CodeGen/Mips/cconv/vector.ll1657
-rw-r--r--test/CodeGen/Mips/ctlz-v.ll12
-rw-r--r--test/CodeGen/Mips/cttz-v.ll19
-rw-r--r--test/CodeGen/Mips/dsp-r1.ll12
-rw-r--r--test/CodeGen/Mips/fmadd1.ll455
-rw-r--r--test/CodeGen/Mips/llvm-ir/mul.ll2
-rw-r--r--test/CodeGen/Mips/llvm-ir/sdiv.ll12
-rw-r--r--test/CodeGen/Mips/llvm-ir/srem.ll11
-rw-r--r--test/CodeGen/Mips/llvm-ir/udiv.ll11
-rw-r--r--test/CodeGen/Mips/llvm-ir/urem.ll6
-rw-r--r--test/CodeGen/Mips/micromips-gp-rc.ll2
-rw-r--r--test/CodeGen/Mips/mips64fpldst.ll12
-rw-r--r--test/CodeGen/Mips/pbqp-reserved-physreg.ll35
-rw-r--r--test/CodeGen/Mips/return-vector.ll33
-rw-r--r--test/CodeGen/Mips/tailcall/tailcall.ll4
-rw-r--r--test/CodeGen/PowerPC/BoolRetToIntTest-2.ll19
-rw-r--r--test/CodeGen/PowerPC/BoolRetToIntTest.ll28
-rw-r--r--test/CodeGen/PowerPC/crbits.ll46
-rw-r--r--test/CodeGen/PowerPC/logic-ops-on-compares.ll73
-rw-r--r--test/CodeGen/PowerPC/memCmpUsedInZeroEqualityComparison.ll277
-rw-r--r--test/CodeGen/PowerPC/ppc-crbits-onoff.ll18
-rw-r--r--test/CodeGen/PowerPC/setcc-logic.ll16
-rw-r--r--test/CodeGen/PowerPC/testComparesinesc.ll121
-rw-r--r--test/CodeGen/PowerPC/testComparesinesi.ll121
-rw-r--r--test/CodeGen/PowerPC/testComparesinesll.ll125
-rw-r--r--test/CodeGen/PowerPC/testComparesiness.ll121
-rw-r--r--test/CodeGen/PowerPC/testComparesineuc.ll136
-rw-r--r--test/CodeGen/PowerPC/testComparesineui.ll121
-rw-r--r--test/CodeGen/PowerPC/testComparesineull.ll125
-rw-r--r--test/CodeGen/PowerPC/testComparesineus.ll137
-rw-r--r--test/CodeGen/PowerPC/testComparesllnesll.ll125
-rw-r--r--test/CodeGen/PowerPC/testComparesllneull.ll125
-rw-r--r--test/CodeGen/PowerPC/vec_int_ext.ll90
-rw-r--r--test/CodeGen/X86/2006-05-11-InstrSched.ll2
-rw-r--r--test/CodeGen/X86/GlobalISel/irtranslator-call.ll30
-rw-r--r--test/CodeGen/X86/GlobalISel/irtranslator-callingconv.ll122
-rw-r--r--test/CodeGen/X86/GlobalISel/irtranslator-callingconv_64bit.ll25
-rw-r--r--test/CodeGen/X86/GlobalISel/legalize-mul-scalar.mir18
-rw-r--r--test/CodeGen/X86/GlobalISel/legalize-mul-v128.mir18
-rw-r--r--test/CodeGen/X86/GlobalISel/legalize-mul-v256.mir18
-rw-r--r--test/CodeGen/X86/GlobalISel/legalize-mul-v512.mir18
-rw-r--r--test/CodeGen/X86/GlobalISel/regbankselect-AVX2.mir20
-rw-r--r--test/CodeGen/X86/GlobalISel/regbankselect-AVX512.mir20
-rw-r--r--test/CodeGen/X86/GlobalISel/regbankselect-X32.mir10
-rw-r--r--test/CodeGen/X86/GlobalISel/regbankselect-X86_64.mir170
-rw-r--r--test/CodeGen/X86/GlobalISel/select-add-v128.mir72
-rw-r--r--test/CodeGen/X86/GlobalISel/select-add-v256.mir72
-rw-r--r--test/CodeGen/X86/GlobalISel/select-add-v512.mir24
-rw-r--r--test/CodeGen/X86/GlobalISel/select-add-x32.mir20
-rw-r--r--test/CodeGen/X86/GlobalISel/select-add.mir72
-rw-r--r--test/CodeGen/X86/GlobalISel/select-cmp.mir130
-rw-r--r--test/CodeGen/X86/GlobalISel/select-constant.mir12
-rw-r--r--test/CodeGen/X86/GlobalISel/select-ext-x86-64.mir16
-rw-r--r--test/CodeGen/X86/GlobalISel/select-ext.mir22
-rw-r--r--test/CodeGen/X86/GlobalISel/select-gep.mir6
-rw-r--r--test/CodeGen/X86/GlobalISel/select-inc.mir8
-rw-r--r--test/CodeGen/X86/GlobalISel/select-leaf-constant.mir8
-rw-r--r--test/CodeGen/X86/GlobalISel/select-memop-scalar-x32.mir56
-rw-r--r--test/CodeGen/X86/GlobalISel/select-memop-scalar.mir76
-rw-r--r--test/CodeGen/X86/GlobalISel/select-memop-v128.mir24
-rw-r--r--test/CodeGen/X86/GlobalISel/select-memop-v256.mir32
-rw-r--r--test/CodeGen/X86/GlobalISel/select-memop-v512.mir16
-rw-r--r--test/CodeGen/X86/GlobalISel/select-mul-scalar.mir18
-rw-r--r--test/CodeGen/X86/GlobalISel/select-mul-vec.mir90
-rw-r--r--test/CodeGen/X86/GlobalISel/select-sub-v128.mir72
-rw-r--r--test/CodeGen/X86/GlobalISel/select-sub-v256.mir72
-rw-r--r--test/CodeGen/X86/GlobalISel/select-sub-v512.mir24
-rw-r--r--test/CodeGen/X86/GlobalISel/select-sub.mir60
-rw-r--r--test/CodeGen/X86/GlobalISel/select-trunc.mir24
-rw-r--r--test/CodeGen/X86/O0-pipeline.ll2
-rw-r--r--test/CodeGen/X86/atom-fixup-lea3.ll11
-rw-r--r--test/CodeGen/X86/avx-schedule.ll32
-rw-r--r--test/CodeGen/X86/avx-splat.ll36
-rw-r--r--test/CodeGen/X86/avx512-cvt.ll4
-rw-r--r--test/CodeGen/X86/build-vector-128.ll92
-rw-r--r--test/CodeGen/X86/buildvec-insertvec.ll72
-rw-r--r--test/CodeGen/X86/clear_upper_vector_element_bits.ll195
-rw-r--r--test/CodeGen/X86/fast-isel-nontemporal.ll108
-rw-r--r--test/CodeGen/X86/full-lsr.ll10
-rw-r--r--test/CodeGen/X86/haddsub-2.ll348
-rw-r--r--test/CodeGen/X86/haddsub-undef.ll5
-rw-r--r--test/CodeGen/X86/hoist-spill.ll2
-rw-r--r--test/CodeGen/X86/loop-strength-reduce4.ll15
-rw-r--r--test/CodeGen/X86/madd.ll78
-rw-r--r--test/CodeGen/X86/masked-iv-safe.ll16
-rw-r--r--test/CodeGen/X86/memcmp.ll330
-rw-r--r--test/CodeGen/X86/merge-consecutive-loads-128.ll36
-rw-r--r--test/CodeGen/X86/mul-constant-i16.ll139
-rw-r--r--test/CodeGen/X86/mul-constant-i32.ll1578
-rw-r--r--test/CodeGen/X86/mul-constant-i64.ll1605
-rw-r--r--test/CodeGen/X86/mul-constant-result.ll1291
-rw-r--r--test/CodeGen/X86/nontemporal-loads.ll779
-rw-r--r--test/CodeGen/X86/pr32659.ll83
-rw-r--r--test/CodeGen/X86/select.ll14
-rw-r--r--test/CodeGen/X86/selectiondag-dominator.ll30
-rw-r--r--test/CodeGen/X86/sse-intrinsics-fast-isel.ll74
-rw-r--r--test/CodeGen/X86/sse1.ll80
-rw-r--r--test/CodeGen/X86/sse2-intrinsics-fast-isel.ll159
-rw-r--r--test/CodeGen/X86/sse3-avx-addsub-2.ll14
-rw-r--r--test/CodeGen/X86/sse42-intrinsics-fast-isel.ll6
-rw-r--r--test/CodeGen/X86/stack-folding-fp-avx1.ll21
-rw-r--r--test/CodeGen/X86/stack-folding-int-sse42.ll17
-rw-r--r--test/CodeGen/X86/trunc-to-bool.ll70
-rw-r--r--test/CodeGen/X86/vec_fp_to_int.ll18
-rw-r--r--test/CodeGen/X86/vec_int_to_fp.ll286
-rw-r--r--test/CodeGen/X86/vec_set.ll24
-rw-r--r--test/CodeGen/X86/vector-compare-results.ll538
-rw-r--r--test/CodeGen/X86/vector-rem.ll34
-rw-r--r--test/CodeGen/X86/vector-sext.ll784
-rw-r--r--test/CodeGen/X86/vector-shuffle-v48.ll49
-rw-r--r--test/CodeGen/X86/vector-shuffle-variable-128.ll274
-rw-r--r--test/CodeGen/X86/vector-sqrt.ll18
-rw-r--r--test/CodeGen/X86/vector-unsigned-cmp.ll134
-rw-r--r--test/CodeGen/X86/virtual-registers-cleared-in-machine-functions-liveins.ll4
-rw-r--r--test/CodeGen/X86/vshift-1.ll9
-rw-r--r--test/CodeGen/X86/vshift-2.ll9
-rw-r--r--test/CodeGen/X86/x86-interleaved-access.ll93
-rw-r--r--test/DebugInfo/Inputs/dwarfdump-str-offsets-dwp.s277
-rw-r--r--test/DebugInfo/Inputs/dwarfdump-str-offsets-invalid-1.s34
-rw-r--r--test/DebugInfo/Inputs/dwarfdump-str-offsets-invalid-1.x86_64.obin0 -> 824 bytes
-rw-r--r--test/DebugInfo/Inputs/dwarfdump-str-offsets-invalid-2.s36
-rw-r--r--test/DebugInfo/Inputs/dwarfdump-str-offsets-invalid-2.x86_64.obin0 -> 832 bytes
-rw-r--r--test/DebugInfo/Inputs/dwarfdump-str-offsets-invalid-3.s88
-rw-r--r--test/DebugInfo/Inputs/dwarfdump-str-offsets-invalid-3.x86_64.obin0 -> 2296 bytes
-rw-r--r--test/DebugInfo/Inputs/dwarfdump-str-offsets-invalid-4.s50
-rw-r--r--test/DebugInfo/Inputs/dwarfdump-str-offsets-invalid-4.x86_64.obin0 -> 1264 bytes
-rw-r--r--test/DebugInfo/Inputs/dwarfdump-str-offsets-invalid-5.s10
-rw-r--r--test/DebugInfo/Inputs/dwarfdump-str-offsets-invalid-5.x86_64.obin0 -> 464 bytes
-rw-r--r--test/DebugInfo/Inputs/dwarfdump-str-offsets.s500
-rw-r--r--test/DebugInfo/Inputs/dwarfdump-str-offsets.x86_64.obin0 -> 4000 bytes
-rw-r--r--test/DebugInfo/PDB/DIA/pdbdump-flags.test8
-rw-r--r--test/DebugInfo/PDB/DIA/pdbdump-linenumbers.test4
-rw-r--r--test/DebugInfo/PDB/DIA/pdbdump-symbol-format.test10
-rw-r--r--test/DebugInfo/PDB/Inputs/debug-subsections.yaml91
-rw-r--r--test/DebugInfo/PDB/Inputs/simple-line-info.yaml44
-rw-r--r--test/DebugInfo/PDB/Native/pdb-native-compilands.test4
-rw-r--r--test/DebugInfo/PDB/Native/pdb-native-summary.test2
-rw-r--r--test/DebugInfo/PDB/pdb-longname-truncation.test2
-rw-r--r--test/DebugInfo/PDB/pdb-minimal-construct.test22
-rw-r--r--test/DebugInfo/PDB/pdb-yaml-symbols.test2
-rw-r--r--test/DebugInfo/PDB/pdb-yaml-types.test6
-rw-r--r--test/DebugInfo/PDB/pdbdump-debug-subsections.test210
-rw-r--r--test/DebugInfo/PDB/pdbdump-headers.test70
-rw-r--r--test/DebugInfo/PDB/pdbdump-merge-ids-and-types.test18
-rw-r--r--test/DebugInfo/PDB/pdbdump-mergeids.test12
-rw-r--r--test/DebugInfo/PDB/pdbdump-mergetypes.test10
-rw-r--r--test/DebugInfo/PDB/pdbdump-raw-blocks.test10
-rw-r--r--test/DebugInfo/PDB/pdbdump-raw-stream.test4
-rw-r--r--test/DebugInfo/PDB/pdbdump-readwrite.test10
-rw-r--r--test/DebugInfo/PDB/pdbdump-source-names.test8
-rw-r--r--test/DebugInfo/PDB/pdbdump-write.test14
-rw-r--r--test/DebugInfo/PDB/pdbdump-yaml-lineinfo-write.test71
-rw-r--r--test/DebugInfo/PDB/pdbdump-yaml-lineinfo.test60
-rw-r--r--test/DebugInfo/PDB/pdbdump-yaml-types.test2
-rw-r--r--test/DebugInfo/PDB/pdbdump-yaml.test4
-rw-r--r--test/DebugInfo/dwarfdump-str-offsets-invalid.test24
-rw-r--r--test/DebugInfo/dwarfdump-str-offsets.test76
-rw-r--r--test/FileCheck/check-dag.txt9
-rw-r--r--test/Instrumentation/MemorySanitizer/csr.ll2
-rw-r--r--test/Instrumentation/MemorySanitizer/msan_x86intrinsics.ll2
-rw-r--r--test/Instrumentation/MemorySanitizer/vector_arith.ll2
-rw-r--r--test/Instrumentation/MemorySanitizer/vector_cmp.ll2
-rw-r--r--test/Instrumentation/MemorySanitizer/vector_cvt.ll2
-rw-r--r--test/Instrumentation/MemorySanitizer/vector_pack.ll2
-rw-r--r--test/Instrumentation/MemorySanitizer/vector_shift.ll2
-rw-r--r--test/LTO/ARM/Inputs/thumb.ll15
-rw-r--r--test/LTO/ARM/link-arm-and-thumb.ll32
-rw-r--r--test/LTO/Resolution/X86/linker-redef.ll16
-rw-r--r--test/Linker/Inputs/thumb.ll16
-rw-r--r--test/Linker/link-arm-and-thumb.ll23
-rw-r--r--test/MC/AMDGPU/sopp-err.s8
-rw-r--r--test/MC/AMDGPU/sym_option.s4
-rw-r--r--test/MC/ARM/arm-thumb-tail-call.ll25
-rw-r--r--test/MC/ARM/big-endian-thumb2-fixup.s6
-rw-r--r--test/MC/ARM/t2-modified-immediate-fixup-error1.s13
-rw-r--r--test/MC/ARM/t2-modified-immediate-fixup-error2.s12
-rw-r--r--test/MC/ARM/t2-modified-immediate-fixup.s45
-rw-r--r--test/MC/ARM/thumb2-diagnostics.s2
-rw-r--r--test/MC/AsmParser/empty-comment.s4
-rw-r--r--test/MC/Disassembler/Mips/micromips-dsp/valid.txt2
-rw-r--r--test/MC/ELF/ARM/clang-section.s399
-rw-r--r--test/MC/MachO/alias.s12
-rw-r--r--test/MC/MachO/variable-exprs.s8
-rw-r--r--test/MC/Mips/dsp/invalid.s4
-rw-r--r--test/MC/Mips/micromips-dsp/invalid.s2
-rw-r--r--test/MC/Mips/micromips-dsp/valid.s2
-rw-r--r--test/MC/WebAssembly/reloc-code.ll20
-rw-r--r--test/Object/AMDGPU/elf-definitions.yaml21
-rw-r--r--test/Object/objc-imageinfo-coff.ll15
-rw-r--r--test/Object/objc-imageinfo-elf.ll15
-rw-r--r--test/Object/objc-imageinfo-macho.ll15
-rw-r--r--test/Transforms/CodeGenPrepare/X86/memcmp.ll337
-rw-r--r--test/Transforms/ConstProp/sse.ll2
-rw-r--r--test/Transforms/DCE/calls-errno.ll4
-rw-r--r--test/Transforms/GVNSink/sink-common-code.ll57
-rw-r--r--test/Transforms/IRCE/correct-loop-info.ll182
-rw-r--r--test/Transforms/IndVarSimplify/lftr_disabled.ll28
-rw-r--r--test/Transforms/InferAddressSpaces/NVPTX/clone_constexpr.ll36
-rw-r--r--test/Transforms/Inline/basictest.ll24
-rw-r--r--test/Transforms/InstCombine/constant-fold-libfunc.ll20
-rw-r--r--test/Transforms/InstCombine/insert-extract-shuffle.ll23
-rw-r--r--test/Transforms/InstCombine/intrinsics.ll154
-rw-r--r--test/Transforms/InstCombine/lshr.ll43
-rw-r--r--test/Transforms/InstSimplify/call.ll33
-rw-r--r--test/Transforms/InstSimplify/compare.ll2
-rw-r--r--test/Transforms/InstSimplify/simplify-nested-bitcast.ll54
-rw-r--r--test/Transforms/InstSimplify/vector_gep.ll2
-rw-r--r--test/Transforms/InterleavedAccess/X86/interleaved-accesses-64bits-avx.ll61
-rw-r--r--test/Transforms/LoopIdiom/X86/unordered-atomic-memcpy.ll452
-rw-r--r--test/Transforms/LoopIdiom/unordered-atomic-memcpy-noarch.ll28
-rw-r--r--test/Transforms/LoopStrengthReduce/X86/canonical.ll2
-rw-r--r--test/Transforms/LoopStrengthReduce/X86/ivchain-X86.ll4
-rw-r--r--test/Transforms/LoopStrengthReduce/X86/lsr-expand-quadratic.ll14
-rw-r--r--test/Transforms/LoopStrengthReduce/X86/lsr-insns-1.ll4
-rw-r--r--test/Transforms/LoopStrengthReduce/X86/lsr-insns-2.ll4
-rw-r--r--test/Transforms/LoopStrengthReduce/X86/nested-loop.ll22
-rw-r--r--test/Transforms/LoopVectorize/AArch64/loop-vectorization-factors.ll46
-rw-r--r--test/Transforms/LowerExpectIntrinsic/PR33346.ll22
-rw-r--r--test/Transforms/LowerTypeTests/simple.ll16
-rw-r--r--test/Transforms/LowerTypeTests/simplify.ll37
-rw-r--r--test/Transforms/NewGVN/completeness.ll2
-rw-r--r--test/Transforms/NewGVN/loadforward.ll4
-rw-r--r--test/Transforms/NewGVN/pr32403.ll3
-rw-r--r--test/Transforms/NewGVN/pr32897.ll1
-rw-r--r--test/Transforms/NewGVN/pr33187.ll148
-rw-r--r--test/Transforms/SLPVectorizer/X86/arith-fp.ll48
-rw-r--r--test/Transforms/SLPVectorizer/X86/reverse_extract_elements.ll138
-rw-r--r--test/Transforms/SROA/address-spaces.ll28
-rw-r--r--test/Transforms/SampleProfile/Inputs/indirect-call.prof3
-rw-r--r--test/Transforms/SampleProfile/indirect-call.ll13
-rw-r--r--test/Transforms/Sink/badloadsink.ll18
-rw-r--r--test/Transforms/ThinLTOBitcodeWriter/split.ll4
-rw-r--r--test/Transforms/Util/PredicateInfo/condprop2.ll2
-rw-r--r--test/Transforms/Util/PredicateInfo/testandor2.ll2
-rw-r--r--test/lit.cfg3
-rw-r--r--test/lit.site.cfg.in1
-rw-r--r--test/tools/llvm-cvtres/Inputs/test_resource.obj.coffbin0 -> 3468 bytes
-rw-r--r--test/tools/llvm-cvtres/object.test229
-rw-r--r--test/tools/llvm-cvtres/parse.test (renamed from test/tools/llvm-cvtres/resource.test)2
-rw-r--r--test/tools/llvm-dwarfdump/X86/brief.s131
-rw-r--r--test/tools/llvm-dwarfdump/X86/lit.local.cfg2
-rw-r--r--test/tools/llvm-pdbdump/class-layout.test2
-rw-r--r--test/tools/llvm-pdbdump/complex-padding-graphical.test2
-rw-r--r--test/tools/llvm-pdbdump/enum-layout.test2
-rw-r--r--test/tools/llvm-pdbdump/load-address.test4
-rw-r--r--test/tools/llvm-pdbdump/raw-stream-data.test6
-rw-r--r--test/tools/llvm-pdbdump/regex-filter.test20
-rw-r--r--test/tools/llvm-pdbdump/simple-padding-graphical.test2
-rw-r--r--test/tools/llvm-pdbdump/symbol-filters.test16
-rwxr-xr-xtest/tools/llvm-readobj/Inputs/trivial.elf-amdhsa-kaveribin13208 -> 0 bytes
-rw-r--r--test/tools/llvm-readobj/Inputs/trivial.obj.elf-amdhsa-gfx803bin0 -> 2208 bytes
-rw-r--r--test/tools/llvm-readobj/amdgpu-elf-definitions.test11
-rw-r--r--test/tools/llvm-readobj/amdgpu-elf-defs.test28
-rw-r--r--test/tools/llvm-readobj/elf-sec-flags.test29
-rw-r--r--tools/LLVMBuild.txt2
-rw-r--r--tools/bugpoint/OptimizerDriver.cpp5
-rw-r--r--tools/dsymutil/DwarfLinker.cpp9
-rw-r--r--tools/llc/llc.cpp97
-rw-r--r--tools/llvm-bcanalyzer/llvm-bcanalyzer.cpp3
-rw-r--r--tools/llvm-cvtres/LLVMBuild.txt2
-rw-r--r--tools/llvm-cvtres/llvm-cvtres.cpp73
-rw-r--r--tools/llvm-cvtres/llvm-cvtres.h2
-rw-r--r--tools/llvm-dwarfdump/llvm-dwarfdump.cpp4
-rw-r--r--tools/llvm-lto2/llvm-lto2.cpp2
-rw-r--r--tools/llvm-mc/llvm-mc.cpp22
-rw-r--r--tools/llvm-nm/llvm-nm.cpp2
-rw-r--r--tools/llvm-objdump/MachODump.cpp7
-rw-r--r--tools/llvm-pdbdump/C13DebugFragmentVisitor.cpp87
-rw-r--r--tools/llvm-pdbdump/C13DebugFragmentVisitor.h60
-rw-r--r--tools/llvm-pdbutil/Analyze.cpp (renamed from tools/llvm-pdbdump/Analyze.cpp)0
-rw-r--r--tools/llvm-pdbutil/Analyze.h (renamed from tools/llvm-pdbdump/Analyze.h)0
-rw-r--r--tools/llvm-pdbutil/CMakeLists.txt (renamed from tools/llvm-pdbdump/CMakeLists.txt)5
-rw-r--r--tools/llvm-pdbutil/CompactTypeDumpVisitor.cpp (renamed from tools/llvm-pdbdump/CompactTypeDumpVisitor.cpp)0
-rw-r--r--tools/llvm-pdbutil/CompactTypeDumpVisitor.h (renamed from tools/llvm-pdbdump/CompactTypeDumpVisitor.h)0
-rw-r--r--tools/llvm-pdbutil/Diff.cpp (renamed from tools/llvm-pdbdump/Diff.cpp)2
-rw-r--r--tools/llvm-pdbutil/Diff.h (renamed from tools/llvm-pdbdump/Diff.h)0
-rw-r--r--tools/llvm-pdbutil/LLVMBuild.txt (renamed from tools/llvm-pdbdump/LLVMBuild.txt)4
-rw-r--r--tools/llvm-pdbutil/LLVMOutputStyle.cpp (renamed from tools/llvm-pdbdump/LLVMOutputStyle.cpp)308
-rw-r--r--tools/llvm-pdbutil/LLVMOutputStyle.h (renamed from tools/llvm-pdbdump/LLVMOutputStyle.h)0
-rw-r--r--tools/llvm-pdbutil/LinePrinter.cpp (renamed from tools/llvm-pdbdump/LinePrinter.cpp)2
-rw-r--r--tools/llvm-pdbutil/LinePrinter.h (renamed from tools/llvm-pdbdump/LinePrinter.h)0
-rw-r--r--tools/llvm-pdbutil/OutputStyle.h (renamed from tools/llvm-pdbdump/OutputStyle.h)0
-rw-r--r--tools/llvm-pdbutil/PdbYaml.cpp (renamed from tools/llvm-pdbdump/PdbYaml.cpp)0
-rw-r--r--tools/llvm-pdbutil/PdbYaml.h (renamed from tools/llvm-pdbdump/PdbYaml.h)0
-rw-r--r--tools/llvm-pdbutil/PrettyBuiltinDumper.cpp (renamed from tools/llvm-pdbdump/PrettyBuiltinDumper.cpp)2
-rw-r--r--tools/llvm-pdbutil/PrettyBuiltinDumper.h (renamed from tools/llvm-pdbdump/PrettyBuiltinDumper.h)0
-rw-r--r--tools/llvm-pdbutil/PrettyClassDefinitionDumper.cpp (renamed from tools/llvm-pdbdump/PrettyClassDefinitionDumper.cpp)2
-rw-r--r--tools/llvm-pdbutil/PrettyClassDefinitionDumper.h (renamed from tools/llvm-pdbdump/PrettyClassDefinitionDumper.h)0
-rw-r--r--tools/llvm-pdbutil/PrettyClassLayoutGraphicalDumper.cpp (renamed from tools/llvm-pdbdump/PrettyClassLayoutGraphicalDumper.cpp)2
-rw-r--r--tools/llvm-pdbutil/PrettyClassLayoutGraphicalDumper.h (renamed from tools/llvm-pdbdump/PrettyClassLayoutGraphicalDumper.h)0
-rw-r--r--tools/llvm-pdbutil/PrettyCompilandDumper.cpp (renamed from tools/llvm-pdbdump/PrettyCompilandDumper.cpp)4
-rw-r--r--tools/llvm-pdbutil/PrettyCompilandDumper.h (renamed from tools/llvm-pdbdump/PrettyCompilandDumper.h)2
-rw-r--r--tools/llvm-pdbutil/PrettyEnumDumper.cpp (renamed from tools/llvm-pdbdump/PrettyEnumDumper.cpp)2
-rw-r--r--tools/llvm-pdbutil/PrettyEnumDumper.h (renamed from tools/llvm-pdbdump/PrettyEnumDumper.h)0
-rw-r--r--tools/llvm-pdbutil/PrettyExternalSymbolDumper.cpp (renamed from tools/llvm-pdbdump/PrettyExternalSymbolDumper.cpp)0
-rw-r--r--tools/llvm-pdbutil/PrettyExternalSymbolDumper.h (renamed from tools/llvm-pdbdump/PrettyExternalSymbolDumper.h)0
-rw-r--r--tools/llvm-pdbutil/PrettyFunctionDumper.cpp (renamed from tools/llvm-pdbdump/PrettyFunctionDumper.cpp)2
-rw-r--r--tools/llvm-pdbutil/PrettyFunctionDumper.h (renamed from tools/llvm-pdbdump/PrettyFunctionDumper.h)0
-rw-r--r--tools/llvm-pdbutil/PrettyTypeDumper.cpp (renamed from tools/llvm-pdbdump/PrettyTypeDumper.cpp)134
-rw-r--r--tools/llvm-pdbutil/PrettyTypeDumper.h (renamed from tools/llvm-pdbdump/PrettyTypeDumper.h)0
-rw-r--r--tools/llvm-pdbutil/PrettyTypedefDumper.cpp (renamed from tools/llvm-pdbdump/PrettyTypedefDumper.cpp)2
-rw-r--r--tools/llvm-pdbutil/PrettyTypedefDumper.h (renamed from tools/llvm-pdbdump/PrettyTypedefDumper.h)2
-rw-r--r--tools/llvm-pdbutil/PrettyVariableDumper.cpp (renamed from tools/llvm-pdbdump/PrettyVariableDumper.cpp)2
-rw-r--r--tools/llvm-pdbutil/PrettyVariableDumper.h (renamed from tools/llvm-pdbdump/PrettyVariableDumper.h)0
-rw-r--r--tools/llvm-pdbutil/StreamUtil.cpp (renamed from tools/llvm-pdbdump/StreamUtil.cpp)0
-rw-r--r--tools/llvm-pdbutil/StreamUtil.h (renamed from tools/llvm-pdbdump/StreamUtil.h)0
-rw-r--r--tools/llvm-pdbutil/YAMLOutputStyle.cpp (renamed from tools/llvm-pdbdump/YAMLOutputStyle.cpp)75
-rw-r--r--tools/llvm-pdbutil/YAMLOutputStyle.h (renamed from tools/llvm-pdbdump/YAMLOutputStyle.h)0
-rw-r--r--tools/llvm-pdbutil/fuzzer/CMakeLists.txt (renamed from tools/llvm-pdbdump/fuzzer/CMakeLists.txt)6
-rw-r--r--tools/llvm-pdbutil/fuzzer/llvm-pdbutil-fuzzer.cpp (renamed from tools/llvm-pdbdump/fuzzer/llvm-pdbdump-fuzzer.cpp)4
-rw-r--r--tools/llvm-pdbutil/llvm-pdbutil.cpp (renamed from tools/llvm-pdbdump/llvm-pdbdump.cpp)140
-rw-r--r--tools/llvm-pdbutil/llvm-pdbutil.h (renamed from tools/llvm-pdbdump/llvm-pdbdump.h)41
-rw-r--r--tools/llvm-readobj/CMakeLists.txt1
-rw-r--r--tools/llvm-readobj/COFFDumper.cpp2
-rw-r--r--tools/llvm-readobj/COFFImportDumper.cpp2
-rw-r--r--tools/llvm-readobj/ELFDumper.cpp202
-rw-r--r--tools/llvm-readobj/LLVMBuild.txt2
-rw-r--r--tools/llvm-stress/llvm-stress.cpp3
-rw-r--r--tools/llvm-xray/xray-extract.cc2
-rw-r--r--tools/obj2yaml/dwarf2yaml.cpp1
-rw-r--r--tools/obj2yaml/macho2yaml.cpp2
-rw-r--r--tools/yaml2obj/yaml2elf.cpp2
-rw-r--r--tools/yaml2obj/yaml2macho.cpp6
-rw-r--r--unittests/ADT/DAGDeltaAlgorithmTest.cpp2
-rw-r--r--unittests/ADT/DeltaAlgorithmTest.cpp2
-rw-r--r--unittests/ADT/DenseMapTest.cpp2
-rw-r--r--unittests/ADT/DepthFirstIteratorTest.cpp2
-rw-r--r--unittests/ADT/FoldingSet.cpp2
-rw-r--r--unittests/ADT/HashingTest.cpp2
-rw-r--r--unittests/ADT/ImmutableMapTest.cpp2
-rw-r--r--unittests/ADT/ImmutableSetTest.cpp2
-rw-r--r--unittests/ADT/IteratorTest.cpp2
-rw-r--r--unittests/ADT/MapVectorTest.cpp2
-rw-r--r--unittests/ADT/OptionalTest.cpp2
-rw-r--r--unittests/ADT/PointerEmbeddedIntTest.cpp2
-rw-r--r--unittests/ADT/PointerIntPairTest.cpp2
-rw-r--r--unittests/ADT/PointerSumTypeTest.cpp2
-rw-r--r--unittests/ADT/PointerUnionTest.cpp2
-rw-r--r--unittests/ADT/PostOrderIteratorTest.cpp2
-rw-r--r--unittests/ADT/RangeAdapterTest.cpp2
-rw-r--r--unittests/ADT/ReverseIterationTest.cpp2
-rw-r--r--unittests/ADT/SCCIteratorTest.cpp2
-rw-r--r--unittests/ADT/SmallPtrSetTest.cpp4
-rw-r--r--unittests/ADT/SmallVectorTest.cpp63
-rw-r--r--unittests/ADT/TripleTest.cpp7
-rw-r--r--unittests/ADT/TwineTest.cpp2
-rw-r--r--unittests/ADT/VariadicFunctionTest.cpp4
-rw-r--r--unittests/Analysis/BranchProbabilityInfoTest.cpp2
-rw-r--r--unittests/Analysis/CFGTest.cpp2
-rw-r--r--unittests/Analysis/CMakeLists.txt1
-rw-r--r--unittests/Analysis/GlobalsModRefTest.cpp55
-rw-r--r--unittests/Analysis/LazyCallGraphTest.cpp2
-rw-r--r--unittests/Analysis/MemorySSA.cpp48
-rw-r--r--unittests/Analysis/ProfileSummaryInfoTest.cpp2
-rw-r--r--unittests/Analysis/ScalarEvolutionTest.cpp1
-rw-r--r--unittests/Analysis/TBAATest.cpp2
-rw-r--r--unittests/Analysis/UnrollAnalyzer.cpp4
-rw-r--r--unittests/BinaryFormat/CMakeLists.txt9
-rw-r--r--unittests/BinaryFormat/DwarfTest.cpp (renamed from unittests/Support/DwarfTest.cpp)5
-rw-r--r--unittests/BinaryFormat/TestFileMagic.cpp128
-rw-r--r--unittests/Bitcode/BitstreamReaderTest.cpp2
-rw-r--r--unittests/Bitcode/BitstreamWriterTest.cpp2
-rw-r--r--unittests/CMakeLists.txt1
-rw-r--r--unittests/CodeGen/DIEHashTest.cpp4
-rw-r--r--unittests/CodeGen/MachineInstrBundleIteratorTest.cpp2
-rw-r--r--unittests/DebugInfo/DWARF/DWARFDebugInfoTest.cpp2
-rw-r--r--unittests/DebugInfo/DWARF/DWARFFormValueTest.cpp2
-rw-r--r--unittests/DebugInfo/DWARF/DwarfGenerator.cpp4
-rw-r--r--unittests/ExecutionEngine/ExecutionEngineTest.cpp2
-rw-r--r--unittests/ExecutionEngine/MCJIT/MCJITCAPITest.cpp2
-rw-r--r--unittests/ExecutionEngine/MCJIT/MCJITMultipleModuleTest.cpp2
-rw-r--r--unittests/ExecutionEngine/MCJIT/MCJITTest.cpp2
-rw-r--r--unittests/ExecutionEngine/Orc/CompileOnDemandLayerTest.cpp2
-rw-r--r--unittests/ExecutionEngine/Orc/IndirectionUtilsTest.cpp2
-rw-r--r--unittests/ExecutionEngine/Orc/LazyEmittingLayerTest.cpp2
-rw-r--r--unittests/ExecutionEngine/Orc/ObjectTransformLayerTest.cpp2
-rw-r--r--unittests/ExecutionEngine/Orc/OrcCAPITest.cpp2
-rw-r--r--unittests/ExecutionEngine/Orc/OrcTestCommon.h4
-rw-r--r--unittests/ExecutionEngine/Orc/QueueChannel.h2
-rw-r--r--unittests/ExecutionEngine/Orc/RTDyldObjectLinkingLayerTest.cpp4
-rw-r--r--unittests/IR/AsmWriterTest.cpp2
-rw-r--r--unittests/IR/ConstantRangeTest.cpp17
-rw-r--r--unittests/IR/ConstantsTest.cpp4
-rw-r--r--unittests/IR/DebugTypeODRUniquingTest.cpp2
-rw-r--r--unittests/IR/DominatorTreeTest.cpp52
-rw-r--r--unittests/IR/IRBuilderTest.cpp2
-rw-r--r--unittests/IR/MetadataTest.cpp2
-rw-r--r--unittests/IR/ModuleTest.cpp2
-rw-r--r--unittests/IR/PassManagerTest.cpp2
-rw-r--r--unittests/IR/PatternMatch.cpp2
-rw-r--r--unittests/IR/UserTest.cpp2
-rw-r--r--unittests/IR/ValueTest.cpp2
-rw-r--r--unittests/IR/VerifierTest.cpp4
-rw-r--r--unittests/Linker/LinkModulesTest.cpp4
-rw-r--r--unittests/MC/DwarfLineTables.cpp2
-rw-r--r--unittests/MC/StringTableBuilderTest.cpp2
-rw-r--r--unittests/MI/LiveIntervalTest.cpp15
-rw-r--r--unittests/ProfileData/CoverageMappingTest.cpp2
-rw-r--r--unittests/ProfileData/InstrProfTest.cpp2
-rw-r--r--unittests/ProfileData/SampleProfTest.cpp2
-rw-r--r--unittests/Support/ARMAttributeParser.cpp2
-rw-r--r--unittests/Support/BinaryStreamTest.cpp9
-rw-r--r--unittests/Support/CMakeLists.txt6
-rw-r--r--unittests/Support/CommandLineTest.cpp4
-rw-r--r--unittests/Support/CompressionTest.cpp2
-rw-r--r--unittests/Support/CrashRecoveryTest.cpp2
-rw-r--r--unittests/Support/DataExtractorTest.cpp2
-rw-r--r--unittests/Support/DynamicLibrary/CMakeLists.txt25
-rw-r--r--unittests/Support/DynamicLibrary/DynamicLibraryTest.cpp53
-rw-r--r--unittests/Support/DynamicLibrary/PipSqueak.cxx34
-rw-r--r--unittests/Support/DynamicLibrary/PipSqueak.h13
-rw-r--r--unittests/Support/EndianStreamTest.cpp2
-rw-r--r--unittests/Support/FileOutputBufferTest.cpp2
-rw-r--r--unittests/Support/FormatVariadicTest.cpp2
-rw-r--r--unittests/Support/LEB128Test.cpp4
-rw-r--r--unittests/Support/MD5Test.cpp2
-rw-r--r--unittests/Support/MathExtrasTest.cpp2
-rw-r--r--unittests/Support/MemoryBufferTest.cpp2
-rw-r--r--unittests/Support/MemoryTest.cpp2
-rw-r--r--unittests/Support/Path.cpp83
-rw-r--r--unittests/Support/ProgramTest.cpp4
-rw-r--r--unittests/Support/SpecialCaseListTest.cpp2
-rw-r--r--unittests/Support/SwapByteOrderTest.cpp2
-rw-r--r--unittests/Support/TarWriterTest.cpp2
-rw-r--r--unittests/Support/TargetParserTest.cpp2
-rw-r--r--unittests/Support/TrigramIndexTest.cpp2
-rw-r--r--unittests/Support/YAMLIOTest.cpp1
-rw-r--r--unittests/Support/YAMLParserTest.cpp2
-rw-r--r--unittests/Support/formatted_raw_ostream_test.cpp2
-rw-r--r--unittests/Support/raw_ostream_test.cpp2
-rw-r--r--unittests/Support/raw_pwrite_stream_test.cpp2
-rw-r--r--unittests/Support/raw_sha1_ostream_test.cpp2
-rw-r--r--unittests/Target/AArch64/InstSizes.cpp24
-rw-r--r--unittests/Transforms/Scalar/LoopPassManagerTest.cpp2
-rw-r--r--unittests/Transforms/Utils/CMakeLists.txt1
-rw-r--r--unittests/Transforms/Utils/FunctionComparator.cpp2
-rw-r--r--unittests/Transforms/Utils/OrderedInstructions.cpp65
-rw-r--r--unittests/Transforms/Utils/ValueMapperTest.cpp2
-rw-r--r--utils/FileCheck/FileCheck.cpp2
-rw-r--r--utils/TableGen/CMakeLists.txt1
-rw-r--r--utils/TableGen/TableGen.cpp6
-rw-r--r--utils/TableGen/TableGenBackends.h1
-rw-r--r--utils/TableGen/X86FoldTablesEmitter.cpp732
-rw-r--r--utils/gdb-scripts/prettyprinters.py36
-rwxr-xr-xutils/git-svn/git-llvm5
-rw-r--r--utils/opt-viewer/optrecord.py2
-rwxr-xr-xutils/release/test-release.sh29
1717 files changed, 39945 insertions, 15255 deletions
diff --git a/CMakeLists.txt b/CMakeLists.txt
index a5b96569f9c6..431785d3dd0c 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -385,6 +385,7 @@ else()
option(LLVM_ENABLE_LOCAL_SUBMODULE_VISIBILITY "Compile with -fmodules-local-submodule-visibility." ON)
endif()
option(LLVM_ENABLE_CXX1Y "Compile with C++1y enabled." OFF)
+option(LLVM_ENABLE_CXX1Z "Compile with C++1z enabled." OFF)
option(LLVM_ENABLE_LIBCXX "Use libc++ if available." OFF)
option(LLVM_ENABLE_LLD "Use lld as C and C++ linker." OFF)
option(LLVM_ENABLE_PEDANTIC "Compile with pedantic enabled." ON)
@@ -853,7 +854,6 @@ if( LLVM_INCLUDE_UTILS )
add_subdirectory(utils/not)
add_subdirectory(utils/llvm-lit)
add_subdirectory(utils/yaml-bench)
- add_subdirectory(utils/unittest)
else()
if ( LLVM_INCLUDE_TESTS )
message(FATAL_ERROR "Including tests when not building utils will not work.
@@ -897,6 +897,10 @@ if( LLVM_INCLUDE_TESTS )
endif()
add_subdirectory(test)
add_subdirectory(unittests)
+ if( LLVM_INCLUDE_UTILS )
+ add_subdirectory(utils/unittest)
+ endif()
+
if (WIN32)
# This utility is used to prevent crashing tests from calling Dr. Watson on
# Windows.
diff --git a/bindings/go/llvm/ir.go b/bindings/go/llvm/ir.go
index fe191beb3813..222097034307 100644
--- a/bindings/go/llvm/ir.go
+++ b/bindings/go/llvm/ir.go
@@ -611,6 +611,12 @@ func (t Type) StructElementTypes() []Type {
}
// Operations on array, pointer, and vector types (sequence types)
+func (t Type) Subtypes() (ret []Type) {
+ ret = make([]Type, C.LLVMGetNumContainedTypes(t.C))
+ C.LLVMGetSubtypes(t.C, llvmTypeRefPtr(&ret[0]))
+ return
+}
+
func ArrayType(elementType Type, elementCount int) (t Type) {
t.C = C.LLVMArrayType(elementType.C, C.unsigned(elementCount))
return
diff --git a/bindings/go/llvm/ir_test.go b/bindings/go/llvm/ir_test.go
index c823615a4293..325ee4890f4c 100644
--- a/bindings/go/llvm/ir_test.go
+++ b/bindings/go/llvm/ir_test.go
@@ -134,3 +134,29 @@ func TestDebugLoc(t *testing.T) {
t.Errorf("Got metadata %v as scope, though wanted %v", loc.Scope.C, scope.C)
}
}
+
+func TestSubtypes(t *testing.T) {
+ cont := NewContext()
+ defer cont.Dispose()
+
+ int_pointer := PointerType(cont.Int32Type(), 0)
+ int_inner := int_pointer.Subtypes()
+ if len(int_inner) != 1 {
+ t.Errorf("Got size %d, though wanted 1")
+ }
+ if int_inner[0] != cont.Int32Type() {
+ t.Errorf("Expected int32 type")
+ }
+
+ st_pointer := cont.StructType([]Type{cont.Int32Type(), cont.Int8Type()}, false)
+ st_inner := st_pointer.Subtypes()
+ if len(st_inner) != 2 {
+ t.Errorf("Got size %d, though wanted 2")
+ }
+ if st_inner[0] != cont.Int32Type() {
+ t.Errorf("Expected first struct field to be int32")
+ }
+ if st_inner[1] != cont.Int8Type() {
+ t.Errorf("Expected second struct field to be int8")
+ }
+}
diff --git a/bindings/ocaml/llvm/llvm.ml b/bindings/ocaml/llvm/llvm.ml
index 399fd2d27c20..6e8ca662ef67 100644
--- a/bindings/ocaml/llvm/llvm.ml
+++ b/bindings/ocaml/llvm/llvm.ml
@@ -459,6 +459,8 @@ external is_packed : lltype -> bool = "llvm_is_packed"
external is_opaque : lltype -> bool = "llvm_is_opaque"
(*--... Operations on pointer, vector, and array types .....................--*)
+
+external subtypes : lltype -> lltype array = "llvm_subtypes"
external array_type : lltype -> int -> lltype = "llvm_array_type"
external pointer_type : lltype -> lltype = "llvm_pointer_type"
external qualified_pointer_type : lltype -> int -> lltype
diff --git a/bindings/ocaml/llvm/llvm.mli b/bindings/ocaml/llvm/llvm.mli
index 4068126e2cbf..c422e78f5d2d 100644
--- a/bindings/ocaml/llvm/llvm.mli
+++ b/bindings/ocaml/llvm/llvm.mli
@@ -658,6 +658,9 @@ val is_opaque : lltype -> bool
(** {7 Operations on pointer, vector, and array types} *)
+(** [subtypes ty] returns [ty]'s subtypes *)
+val subtypes : lltype -> lltype array
+
(** [array_type ty n] returns the array type containing [n] elements of type
[ty]. See the method [llvm::ArrayType::get]. *)
val array_type : lltype -> int -> lltype
diff --git a/bindings/ocaml/llvm/llvm_ocaml.c b/bindings/ocaml/llvm/llvm_ocaml.c
index af04ea25c8ab..4b6d1c5072bc 100644
--- a/bindings/ocaml/llvm/llvm_ocaml.c
+++ b/bindings/ocaml/llvm/llvm_ocaml.c
@@ -506,6 +506,20 @@ CAMLprim value llvm_is_opaque(LLVMTypeRef StructTy) {
/*--... Operations on array, pointer, and vector types .....................--*/
+/* lltype -> lltype array */
+CAMLprim value llvm_subtypes(LLVMTypeRef Ty) {
+ CAMLparam0();
+ CAMLlocal1(Arr);
+
+ unsigned Size = LLVMGetNumContainedTypes(Ty);
+
+ Arr = caml_alloc(Size, 0);
+
+ LLVMGetSubtypes(Ty, (LLVMTypeRef *) Arr);
+
+ CAMLreturn(Arr);
+}
+
/* lltype -> int -> lltype */
CAMLprim LLVMTypeRef llvm_array_type(LLVMTypeRef ElementTy, value Count) {
return LLVMArrayType(ElementTy, Int_val(Count));
diff --git a/cmake/modules/HandleLLVMOptions.cmake b/cmake/modules/HandleLLVMOptions.cmake
index 3dd16d51f0b7..c3325db11788 100644
--- a/cmake/modules/HandleLLVMOptions.cmake
+++ b/cmake/modules/HandleLLVMOptions.cmake
@@ -101,6 +101,10 @@ else()
message(FATAL_ERROR "Unknown value for LLVM_ABI_BREAKING_CHECKS: \"${LLVM_ABI_BREAKING_CHECKS}\"!")
endif()
+if( LLVM_REVERSE_ITERATION )
+ set( LLVM_ENABLE_REVERSE_ITERATION 1 )
+endif()
+
if(WIN32)
set(LLVM_HAVE_LINK_VERSION_SCRIPT 0)
if(CYGWIN)
@@ -381,6 +385,9 @@ elseif( LLVM_COMPILER_IS_GCC_COMPATIBLE )
if (LLVM_ENABLE_CXX1Y)
check_cxx_compiler_flag("-std=c++1y" CXX_SUPPORTS_CXX1Y)
append_if(CXX_SUPPORTS_CXX1Y "-std=c++1y" CMAKE_CXX_FLAGS)
+ elseif(LLVM_ENABLE_CXX1Z)
+ check_cxx_compiler_flag("-std=c++1z" CXX_SUPPORTS_CXX1Z)
+ append_if(CXX_SUPPORTS_CXX1Z "-std=c++1z" CMAKE_CXX_FLAGS)
else()
check_cxx_compiler_flag("-std=c++11" CXX_SUPPORTS_CXX11)
if (CXX_SUPPORTS_CXX11)
diff --git a/cmake/modules/TableGen.cmake b/cmake/modules/TableGen.cmake
index da0858e54d44..17ae1c9e7717 100644
--- a/cmake/modules/TableGen.cmake
+++ b/cmake/modules/TableGen.cmake
@@ -30,19 +30,43 @@ function(tablegen project ofn)
endif()
endif()
- add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/${ofn}.tmp
- # Generate tablegen output in a temporary file.
- COMMAND ${${project}_TABLEGEN_EXE} ${ARGN} -I ${CMAKE_CURRENT_SOURCE_DIR}
- ${LLVM_TABLEGEN_FLAGS}
- ${LLVM_TARGET_DEFINITIONS_ABSOLUTE}
- -o ${CMAKE_CURRENT_BINARY_DIR}/${ofn}.tmp
- # The file in LLVM_TARGET_DEFINITIONS may be not in the current
- # directory and local_tds may not contain it, so we must
- # explicitly list it here:
- DEPENDS ${${project}_TABLEGEN_TARGET} ${local_tds} ${global_tds}
- ${LLVM_TARGET_DEFINITIONS_ABSOLUTE}
- COMMENT "Building ${ofn}..."
- )
+ # We need both _TABLEGEN_TARGET and _TABLEGEN_EXE in the DEPENDS list
+ # (both the target and the file) to have .inc files rebuilt on
+ # a tablegen change, as cmake does not propagate file-level dependencies
+ # of custom targets. See the following ticket for more information:
+ # https://cmake.org/Bug/view.php?id=15858
+ # We could always have just one dependency on both the target and
+ # the file, but these 2 cases would produce cleaner cmake files.
+ if (${${project}_TABLEGEN_TARGET} STREQUAL ${${project}_TABLEGEN_EXE})
+ add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/${ofn}.tmp
+ # Generate tablegen output in a temporary file.
+ COMMAND ${${project}_TABLEGEN_EXE} ${ARGN} -I ${CMAKE_CURRENT_SOURCE_DIR}
+ ${LLVM_TABLEGEN_FLAGS}
+ ${LLVM_TARGET_DEFINITIONS_ABSOLUTE}
+ -o ${CMAKE_CURRENT_BINARY_DIR}/${ofn}.tmp
+ # The file in LLVM_TARGET_DEFINITIONS may be not in the current
+ # directory and local_tds may not contain it, so we must
+ # explicitly list it here:
+ DEPENDS ${${project}_TABLEGEN_TARGET} ${local_tds} ${global_tds}
+ ${LLVM_TARGET_DEFINITIONS_ABSOLUTE}
+ COMMENT "Building ${ofn}..."
+ )
+ else()
+ add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/${ofn}.tmp
+ # Generate tablegen output in a temporary file.
+ COMMAND ${${project}_TABLEGEN_EXE} ${ARGN} -I ${CMAKE_CURRENT_SOURCE_DIR}
+ ${LLVM_TABLEGEN_FLAGS}
+ ${LLVM_TARGET_DEFINITIONS_ABSOLUTE}
+ -o ${CMAKE_CURRENT_BINARY_DIR}/${ofn}.tmp
+ # The file in LLVM_TARGET_DEFINITIONS may be not in the current
+ # directory and local_tds may not contain it, so we must
+ # explicitly list it here:
+ DEPENDS ${${project}_TABLEGEN_TARGET} ${${project}_TABLEGEN_EXE}
+ ${local_tds} ${global_tds}
+ ${LLVM_TARGET_DEFINITIONS_ABSOLUTE}
+ COMMENT "Building ${ofn}..."
+ )
+ endif()
add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/${ofn}
# Only update the real output file if there are any differences.
# This prevents recompilation of all the files depending on it if there
diff --git a/docs/AMDGPUUsage.rst b/docs/AMDGPUUsage.rst
index 81c067b317d3..caa697ca28cd 100644
--- a/docs/AMDGPUUsage.rst
+++ b/docs/AMDGPUUsage.rst
@@ -1,109 +1,3441 @@
-==============================
-User Guide for AMDGPU Back-end
-==============================
+=============================
+User Guide for AMDGPU Backend
+=============================
+
+.. contents::
+ :local:
Introduction
============
-The AMDGPU back-end provides ISA code generation for AMD GPUs, starting with
-the R600 family up until the current Volcanic Islands (GCN Gen 3).
+The AMDGPU backend provides ISA code generation for AMD GPUs, starting with the
+R600 family up until the current GCN families. It lives in the
+``lib/Target/AMDGPU`` directory.
-Refer to `AMDGPU section in Architecture & Platform Information for Compiler Writers <CompilerWriterInfo.html#amdgpu>`_
-for additional documentation.
+LLVM
+====
-Conventions
-===========
+.. _amdgpu-target-triples:
+
+Target Triples
+--------------
+
+Use the ``clang -target <Architecture>-<Vendor>-<OS>-<Environment>`` option to
+specify the target triple:
+
+ .. table:: AMDGPU Target Triples
+ :name: amdgpu-target-triples-table
+
+ ============ ======== ========= ===========
+ Architecture Vendor OS Environment
+ ============ ======== ========= ===========
+ r600 amd <empty> <empty>
+ amdgcn amd <empty> <empty>
+ amdgcn amd amdhsa <empty>
+ amdgcn amd amdhsa opencl
+ amdgcn amd amdhsa amdgizcl
+ amdgcn amd amdhsa amdgiz
+ amdgcn amd amdhsa hcc
+ ============ ======== ========= ===========
+
+``r600-amd--``
+ Supports AMD GPUs HD2XXX-HD6XXX for graphics and compute shaders executed on
+ the MESA runtime.
+
+``amdgcn-amd--``
+ Supports AMD GPUs GCN 6 onwards for graphics and compute shaders executed on
+ the MESA runtime.
+
+``amdgcn-amd-amdhsa-``
+ Supports AMD GCN GPUs GFX6 onwards for compute kernels executed on HSA [HSA]_
+ compatible runtimes such as AMD's ROCm [AMD-ROCm]_.
+
+``amdgcn-amd-amdhsa-opencl``
+ Supports AMD GCN GPUs GFX6 onwards for OpenCL compute kernels executed on HSA
+ [HSA]_ compatible runtimes such as AMD's ROCm [AMD-ROCm]_. See
+ :ref:`amdgpu-opencl`.
+
+``amdgcn-amd-amdhsa-amdgizcl``
+ Same as ``amdgcn-amd-amdhsa-opencl`` except a different address space mapping
+ is used (see :ref:`amdgpu-address-spaces`).
+
+``amdgcn-amd-amdhsa-amdgiz``
+ Same as ``amdgcn-amd-amdhsa-`` except a different address space mapping is
+ used (see :ref:`amdgpu-address-spaces`).
+
+``amdgcn-amd-amdhsa-hcc``
+ Supports AMD GCN GPUs GFX6 onwards for AMD HC language compute kernels
+ executed on HSA [HSA]_ compatible runtimes such as AMD's ROCm [AMD-ROCm]_. See
+ :ref:`amdgpu-hcc`.
+
+.. _amdgpu-processors:
+
+Processors
+----------
+
+Use the ``clang -mcpu <Processor>`` option to specify the AMD GPU processor. The
+names from both the *Processor* and *Alternative Processor* can be used.
+
+ .. table:: AMDGPU Processors
+ :name: amdgpu-processors-table
+
+ ========== =========== ============ ===== ======= ==================
+ Processor Alternative Target dGPU/ Runtime Example
+ Processor Triple APU Support Products
+ Architecture
+ ========== =========== ============ ===== ======= ==================
+ **R600** [AMD-R6xx]_
+ --------------------------------------------------------------------
+ r600 r600 dGPU
+ r630 r600 dGPU
+ rs880 r600 dGPU
+ rv670 r600 dGPU
+ **R700** [AMD-R7xx]_
+ --------------------------------------------------------------------
+ rv710 r600 dGPU
+ rv730 r600 dGPU
+ rv770 r600 dGPU
+ **Evergreen** [AMD-Evergreen]_
+ --------------------------------------------------------------------
+ cedar r600 dGPU
+ redwood r600 dGPU
+ sumo r600 dGPU
+ juniper r600 dGPU
+ cypress r600 dGPU
+ **Northern Islands** [AMD-Cayman-Trinity]_
+ --------------------------------------------------------------------
+ barts r600 dGPU
+ turks r600 dGPU
+ caicos r600 dGPU
+ cayman r600 dGPU
+ **GCN GFX6 (Southern Islands (SI))** [AMD-Souther-Islands]_
+ --------------------------------------------------------------------
+ gfx600 - SI amdgcn dGPU
+ - tahiti
+ gfx601 - pitcairn amdgcn dGPU
+ - verde
+ - oland
+ - hainan
+ **GCN GFX7 (Sea Islands (CI))** [AMD-Sea-Islands]_
+ --------------------------------------------------------------------
+ gfx700 - bonaire amdgcn dGPU - Radeon HD 7790
+ - Radeon HD 8770
+ - R7 260
+ - R7 260X
+ \ - kaveri amdgcn APU - A6-7000
+ - A6 Pro-7050B
+ - A8-7100
+ - A8 Pro-7150B
+ - A10-7300
+ - A10 Pro-7350B
+ - FX-7500
+ - A8-7200P
+ - A10-7400P
+ - FX-7600P
+ gfx701 - hawaii amdgcn dGPU ROCm - FirePro W8100
+ - FirePro W9100
+ - FirePro S9150
+ - FirePro S9170
+ gfx702 dGPU ROCm - Radeon R9 290
+ - Radeon R9 290x
+ - Radeon R390
+ - Radeon R390x
+ gfx703 - kabini amdgcn APU - E1-2100
+ - mullins - E1-2200
+ - E1-2500
+ - E2-3000
+ - E2-3800
+ - A4-5000
+ - A4-5100
+ - A6-5200
+ - A4 Pro-3340B
+ **GCN GFX8 (Volcanic Islands (VI))** [AMD-Volcanic-Islands]_
+ --------------------------------------------------------------------
+ gfx800 - iceland amdgcn dGPU - FirePro S7150
+ - FirePro S7100
+ - FirePro W7100
+ - Radeon R285
+ - Radeon R9 380
+ - Radeon R9 385
+ - Mobile FirePro
+ M7170
+ gfx801 - carrizo amdgcn APU - A6-8500P
+ - Pro A6-8500B
+ - A8-8600P
+ - Pro A8-8600B
+ - FX-8800P
+ - Pro A12-8800B
+ \ amdgcn APU ROCm - A10-8700P
+ - Pro A10-8700B
+ - A10-8780P
+ \ amdgcn APU - A10-9600P
+ - A10-9630P
+ - A12-9700P
+ - A12-9730P
+ - FX-9800P
+ - FX-9830P
+ \ amdgcn APU - E2-9010
+ - A6-9210
+ - A9-9410
+ gfx802 - tonga amdgcn dGPU ROCm Same as gfx800
+ gfx803 - fiji amdgcn dGPU ROCm - Radeon R9 Nano
+ - Radeon R9 Fury
+ - Radeon R9 FuryX
+ - Radeon Pro Duo
+ - FirePro S9300x2
+ \ - polaris10 amdgcn dGPU ROCm - Radeon RX 470
+ - Radeon RX 480
+ \ - polaris11 amdgcn dGPU ROCm - Radeon RX 460
+ gfx804 amdgcn dGPU Same as gfx803
+ gfx810 - stoney amdgcn APU
+ **GCN GFX9**
+ --------------------------------------------------------------------
+ gfx900 amdgcn dGPU - FirePro W9500
+ - FirePro S9500
+ - FirePro S9500x2
+ gfx901 amdgcn dGPU ROCm Same as gfx900
+ except XNACK is
+ enabled
+ gfx902 amdgcn APU *TBA*
+
+ .. TODO
+ Add product
+ names.
+ gfx903 amdgcn APU Same as gfx902
+ except XNACK is
+ enabled
+ ========== =========== ============ ===== ======= ==================
+
+.. _amdgpu-address-spaces:
Address Spaces
--------------
-The AMDGPU back-end uses the following address space mapping:
+The AMDGPU backend uses the following address space mappings.
+
+The memory space names used in the table, aside from the region memory space, is
+from the OpenCL standard.
+
+LLVM Address Space number is used throughout LLVM (for example, in LLVM IR).
+
+ .. table:: Address Space Mapping
+ :name: amdgpu-address-space-mapping-table
+
+ ================== ================= ================= ================= =================
+ LLVM Address Space Memory Space
+ ------------------ -----------------------------------------------------------------------
+ \ Current Default amdgiz/amdgizcl hcc Future Default
+ ================== ================= ================= ================= =================
+ 0 Private (Scratch) Generic (Flat) Generic (Flat) Generic (Flat)
+ 1 Global Global Global Global
+ 2 Constant Constant Constant Region (GDS)
+ 3 Local (group/LDS) Local (group/LDS) Local (group/LDS) Local (group/LDS)
+ 4 Generic (Flat) Region (GDS) Region (GDS) Constant
+ 5 Region (GDS) Private (Scratch) Private (Scratch) Private (Scratch)
+ ================== ================= ================= ================= =================
+
+Current Default
+ This is the current default address space mapping used for all languages
+ except hcc. This will shortly be deprecated.
+
+amdgiz/amdgizcl
+ This is the current address space mapping used when ``amdgiz`` or ``amdgizcl``
+ is specified as the target triple environment value.
+
+hcc
+ This is the current address space mapping used when ``hcc`` is specified as
+ the target triple environment value.This will shortly be deprecated.
+
+Future Default
+ This will shortly be the only address space mapping for all languages using
+ AMDGPU backend.
+
+.. _amdgpu-memory-scopes:
+
+Memory Scopes
+-------------
+
+This section provides LLVM memory synchronization scopes supported by the AMDGPU
+backend memory model when the target triple OS is ``amdhsa`` (see
+:ref:`amdgpu-amdhsa-memory-model` and :ref:`amdgpu-target-triples`).
+
+The memory model supported is based on the HSA memory model [HSA]_ which is
+based in turn on HRF-indirect with scope inclusion [HRF]_. The happens-before
+relation is transitive over the synchonizes-with relation independent of scope,
+and synchonizes-with allows the memory scope instances to be inclusive (see
+table :ref:`amdgpu-amdhsa-llvm-sync-scopes-amdhsa-table`).
+
+This is different to the OpenCL [OpenCL]_ memory model which does not have scope
+inclusion and requires the memory scopes to exactly match. However, this
+is conservatively correct for OpenCL.
+
+ .. table:: AMDHSA LLVM Sync Scopes for AMDHSA
+ :name: amdgpu-amdhsa-llvm-sync-scopes-amdhsa-table
+
+ ================ ==========================================================
+ LLVM Sync Scope Description
+ ================ ==========================================================
+ *none* The default: ``system``.
+
+ Synchronizes with, and participates in modification and
+ seq_cst total orderings with, other operations (except
+ image operations) for all address spaces (except private,
+ or generic that accesses private) provided the other
+ operation's sync scope is:
+
+ - ``system``.
+ - ``agent`` and executed by a thread on the same agent.
+ - ``workgroup`` and executed by a thread in the same
+ workgroup.
+ - ``wavefront`` and executed by a thread in the same
+ wavefront.
+
+ ``agent`` Synchronizes with, and participates in modification and
+ seq_cst total orderings with, other operations (except
+ image operations) for all address spaces (except private,
+ or generic that accesses private) provided the other
+ operation's sync scope is:
+
+ - ``system`` or ``agent`` and executed by a thread on the
+ same agent.
+ - ``workgroup`` and executed by a thread in the same
+ workgroup.
+ - ``wavefront`` and executed by a thread in the same
+ wavefront.
+
+ ``workgroup`` Synchronizes with, and participates in modification and
+ seq_cst total orderings with, other operations (except
+ image operations) for all address spaces (except private,
+ or generic that accesses private) provided the other
+ operation's sync scope is:
+
+ - ``system``, ``agent`` or ``workgroup`` and executed by a
+ thread in the same workgroup.
+ - ``wavefront`` and executed by a thread in the same
+ wavefront.
+
+ ``wavefront`` Synchronizes with, and participates in modification and
+ seq_cst total orderings with, other operations (except
+ image operations) for all address spaces (except private,
+ or generic that accesses private) provided the other
+ operation's sync scope is:
+
+ - ``system``, ``agent``, ``workgroup`` or ``wavefront``
+ and executed by a thread in the same wavefront.
+
+ ``singlethread`` Only synchronizes with, and participates in modification
+ and seq_cst total orderings with, other operations (except
+ image operations) running in the same thread for all
+ address spaces (for example, in signal handlers).
+ ================ ==========================================================
+
+AMDGPU Intrinsics
+-----------------
+
+The AMDGPU backend implements the following intrinsics.
+
+*This section is WIP.*
+
+.. TODO
+ List AMDGPU intrinsics
+
+Code Object
+===========
+
+The AMDGPU backend generates a standard ELF [ELF]_ relocatable code object that
+can be linked by ``lld`` to produce a standard ELF shared code object which can
+be loaded and executed on an AMDGPU target.
+
+Header
+------
+
+The AMDGPU backend uses the following ELF header:
+
+ .. table:: AMDGPU ELF Header
+ :name: amdgpu-elf-header-table
+
+ ========================== =========================
+ Field Value
+ ========================== =========================
+ ``e_ident[EI_CLASS]`` ``ELFCLASS64``
+ ``e_ident[EI_DATA]`` ``ELFDATA2LSB``
+ ``e_ident[EI_OSABI]`` ``ELFOSABI_AMDGPU_HSA``
+ ``e_ident[EI_ABIVERSION]`` ``ELFABIVERSION_AMDGPU_HSA``
+ ``e_type`` ``ET_REL`` or ``ET_DYN``
+ ``e_machine`` ``EM_AMDGPU``
+ ``e_entry`` 0
+ ``e_flags`` 0
+ ========================== =========================
+
+..
+
+ .. table:: AMDGPU ELF Header Enumeration Values
+ :name: amdgpu-elf-header-enumeration-values-table
+
+ ============================ =====
+ Name Value
+ ============================ =====
+ ``EM_AMDGPU`` 224
+ ``ELFOSABI_AMDGPU_HSA`` 64
+ ``ELFABIVERSION_AMDGPU_HSA`` 1
+ ============================ =====
+
+``e_ident[EI_CLASS]``
+ The ELF class is always ``ELFCLASS64``. The AMDGPU backend only supports 64 bit
+ applications.
+
+``e_ident[EI_DATA]``
+ All AMDGPU targets use ELFDATA2LSB for little-endian byte ordering.
+
+``e_ident[EI_OSABI]``
+ The AMD GPU architecture specific OS ABI of ``ELFOSABI_AMDGPU_HSA`` is used to
+ specify that the code object conforms to the AMD HSA runtime ABI [HSA]_.
+
+``e_ident[EI_ABIVERSION]``
+ The AMD GPU architecture specific OS ABI version of
+ ``ELFABIVERSION_AMDGPU_HSA`` is used to specify the version of AMD HSA runtime
+ ABI to which the code object conforms.
+
+``e_type``
+ Can be one of the following values:
+
+
+ ``ET_REL``
+ The type produced by the AMD GPU backend compiler as it is relocatable code
+ object.
+
+ ``ET_DYN``
+ The type produced by the linker as it is a shared code object.
+
+ The AMD HSA runtime loader requires a ``ET_DYN`` code object.
+
+``e_machine``
+ The value ``EM_AMDGPU`` is used for the machine for all members of the AMD GPU
+ architecture family. The specific member is specified in the
+ ``NT_AMD_AMDGPU_ISA`` entry in the ``.note`` section (see
+ :ref:`amdgpu-note-records`).
+
+``e_entry``
+ The entry point is 0 as the entry points for individual kernels must be
+ selected in order to invoke them through AQL packets.
+
+``e_flags``
+ The value is 0 as no flags are used.
+
+Sections
+--------
+
+An AMDGPU target ELF code object has the standard ELF sections which include:
+
+ .. table:: AMDGPU ELF Sections
+ :name: amdgpu-elf-sections-table
+
+ ================== ================ =================================
+ Name Type Attributes
+ ================== ================ =================================
+ ``.bss`` ``SHT_NOBITS`` ``SHF_ALLOC`` + ``SHF_WRITE``
+ ``.data`` ``SHT_PROGBITS`` ``SHF_ALLOC`` + ``SHF_WRITE``
+ ``.debug_``\ *\** ``SHT_PROGBITS`` *none*
+ ``.dynamic`` ``SHT_DYNAMIC`` ``SHF_ALLOC``
+ ``.dynstr`` ``SHT_PROGBITS`` ``SHF_ALLOC``
+ ``.dynsym`` ``SHT_PROGBITS`` ``SHF_ALLOC``
+ ``.got`` ``SHT_PROGBITS`` ``SHF_ALLOC`` + ``SHF_WRITE``
+ ``.hash`` ``SHT_HASH`` ``SHF_ALLOC``
+ ``.note`` ``SHT_NOTE`` *none*
+ ``.rela``\ *name* ``SHT_RELA`` *none*
+ ``.rela.dyn`` ``SHT_RELA`` *none*
+ ``.rodata`` ``SHT_PROGBITS`` ``SHF_ALLOC``
+ ``.shstrtab`` ``SHT_STRTAB`` *none*
+ ``.strtab`` ``SHT_STRTAB`` *none*
+ ``.symtab`` ``SHT_SYMTAB`` *none*
+ ``.text`` ``SHT_PROGBITS`` ``SHF_ALLOC`` + ``SHF_EXECINSTR``
+ ================== ================ =================================
+
+These sections have their standard meanings (see [ELF]_) and are only generated
+if needed.
+
+``.debug``\ *\**
+ The standard DWARF sections. See :ref:`amdgpu-dwarf` for information on the
+ DWARF produced by the AMDGPU backend.
+
+``.dynamic``, ``.dynstr``, ``.dynstr``, ``.hash``
+ The standard sections used by a dynamic loader.
+
+``.note``
+ See :ref:`amdgpu-note-records` for the note records supported by the AMDGPU
+ backend.
+
+``.rela``\ *name*, ``.rela.dyn``
+ For relocatable code objects, *name* is the name of the section that the
+ relocation records apply. For example, ``.rela.text`` is the section name for
+ relocation records associated with the ``.text`` section.
+
+ For linked shared code objects, ``.rela.dyn`` contains all the relocation
+ records from each of the relocatable code object's ``.rela``\ *name* sections.
+
+ See :ref:`amdgpu-relocation-records` for the relocation records supported by
+ the AMDGPU backend.
+
+``.text``
+ The executable machine code for the kernels and functions they call. Generated
+ as position independent code. See :ref:`amdgpu-code-conventions` for
+ information on conventions used in the isa generation.
+
+.. _amdgpu-note-records:
+
+Note Records
+------------
+
+As required by ``ELFCLASS64``, minimal zero byte padding must be generated after
+the ``name`` field to ensure the ``desc`` field is 4 byte aligned. In addition,
+minimal zero byte padding must be generated to ensure the ``desc`` field size is
+a multiple of 4 bytes. The ``sh_addralign`` field of the ``.note`` section must
+be at least 4 to indicate at least 8 byte alignment.
+
+The AMDGPU backend code object uses the following ELF note records in the
+``.note`` section. The *Description* column specifies the layout of the note
+record’s ``desc`` field. All fields are consecutive bytes. Note records with
+variable size strings have a corresponding ``*_size`` field that specifies the
+number of bytes, including the terminating null character, in the string. The
+string(s) come immediately after the preceding fields.
+
+Additional note records can be present.
+
+ .. table:: AMDGPU ELF Note Records
+ :name: amdgpu-elf-note-records-table
- ================== =================== ==============
- LLVM Address Space DWARF Address Space Memory Space
- ================== =================== ==============
- 0 1 Private
- 1 N/A Global
- 2 N/A Constant
- 3 2 Local
- 4 N/A Generic (Flat)
- 5 N/A Region
- ================== =================== ==============
+ ===== ========================== ==========================================
+ Name Type Description
+ ===== ========================== ==========================================
+ "AMD" ``NT_AMD_AMDGPU_METADATA`` <metadata null terminated string>
+ "AMD" ``NT_AMD_AMDGPU_ISA`` <isa name null terminated string>
+ ===== ========================== ==========================================
-The terminology in the table, aside from the region memory space, is from the
-OpenCL standard.
+..
-LLVM Address Space is used throughout LLVM (for example, in LLVM IR). DWARF
-Address Space is emitted in DWARF, and is used by tools, such as debugger,
-profiler and others.
+ .. table:: AMDGPU ELF Note Record Enumeration Values
+ :name: amdgpu-elf-note-record-enumeration-values-table
+
+ ============================= =====
+ Name Value
+ ============================= =====
+ *reserved* 0-9
+ ``NT_AMD_AMDGPU_METADATA`` 10
+ ``NT_AMD_AMDGPU_ISA`` 11
+ ============================= =====
+
+``NT_AMD_AMDGPU_ISA``
+ Specifies the instruction set architecture used by the machine code contained
+ in the code object.
+
+ This note record is required for code objects containing machine code for
+ processors matching the ``amdgcn`` architecture in table
+ :ref:`amdgpu-processors`.
+
+ The null terminated string has the following syntax:
+
+ *architecture*\ ``-``\ *vendor*\ ``-``\ *os*\ ``-``\ *environment*\ ``-``\ *processor*
+
+ where:
+
+ *architecture*
+ The architecture from table :ref:`amdgpu-target-triples-table`.
+
+ This is always ``amdgcn`` when the target triple OS is ``amdhsa`` (see
+ :ref:`amdgpu-target-triples`).
+
+ *vendor*
+ The vendor from table :ref:`amdgpu-target-triples-table`.
+
+ For the AMDGPU backend this is always ``amd``.
+
+ *os*
+ The OS from table :ref:`amdgpu-target-triples-table`.
+
+ *environment*
+ An environment from table :ref:`amdgpu-target-triples-table`, or blank if
+ the environment has no affect on the execution of the code object.
+
+ For the AMDGPU backend this is currently always blank.
+ *processor*
+ The processor from table :ref:`amdgpu-processors-table`.
+
+ For example:
+
+ ``amdgcn-amd-amdhsa--gfx901``
+
+``NT_AMD_AMDGPU_METADATA``
+ Specifies extensible metadata associated with the code object. See
+ :ref:`amdgpu-code-object-metadata` for the syntax of the code object metadata
+ string.
+
+ This note record is required and must contain the minimum information
+ necessary to support the ROCM kernel queries. For example, the segment sizes
+ needed in a dispatch packet. In addition, a high level language runtime may
+ require other information to be included. For example, the AMD OpenCL runtime
+ records kernel argument information.
+
+ .. TODO
+ Is the string null terminated? It probably should not if YAML allows it to
+ contain null characters, otherwise it should be.
+
+.. _amdgpu-code-object-metadata:
+
+Code Object Metadata
+--------------------
+
+The code object metadata is specified by the ``NT_AMD_AMDHSA_METADATA`` note
+record (see :ref:`amdgpu-note-records`).
+
+The metadata is specified as a YAML formated string (see [YAML]_ and
+:doc:`YamlIO`).
+
+The metadata is represented as a single YAML document comprised of the mapping
+defined in table :ref:`amdgpu-amdhsa-code-object-metadata-mapping-table` and
+referenced tables.
+
+For boolean values, the string values of ``false`` and ``true`` are used for
+false and true respectively.
+
+Additional information can be added to the mappings. To avoid conflicts, any
+non-AMD key names should be prefixed by "*vendor-name*.".
+
+ .. table:: AMDHSA Code Object Metadata Mapping
+ :name: amdgpu-amdhsa-code-object-metadata-mapping-table
+
+ ========== ============== ========= =======================================
+ String Key Value Type Required? Description
+ ========== ============== ========= =======================================
+ "Version" sequence of Required - The first integer is the major
+ 2 integers version. Currently 1.
+ - The second integer is the minor
+ version. Currently 0.
+ "Printf" sequence of Each string is encoded information
+ strings about a printf function call. The
+ encoded information is organized as
+ fields separated by colon (':'):
+
+ ``ID:N:S[0]:S[1]:...:S[N-1]:FormatString``
+
+ where:
+
+ ``ID``
+ A 32 bit integer as a unique id for
+ each printf function call
+
+ ``N``
+ A 32 bit integer equal to the number
+ of arguments of printf function call
+ minus 1
+
+ ``S[i]`` (where i = 0, 1, ... , N-1)
+ 32 bit integers for the size in bytes
+ of the i-th FormatString argument of
+ the printf function call
+
+ FormatString
+ The format string passed to the
+ printf function call.
+ "Kernels" sequence of Required Sequence of the mappings for each
+ mapping kernel in the code object. See
+ :ref:`amdgpu-amdhsa-code-object-kernel-metadata-mapping-table`
+ for the definition of the mapping.
+ ========== ============== ========= =======================================
+
+..
+
+ .. table:: AMDHSA Code Object Kernel Metadata Mapping
+ :name: amdgpu-amdhsa-code-object-kernel-metadata-mapping-table
+
+ ================= ============== ========= ================================
+ String Key Value Type Required? Description
+ ================= ============== ========= ================================
+ "Name" string Required Source name of the kernel.
+ "SymbolName" string Required Name of the kernel
+ descriptor ELF symbol.
+ "Language" string Source language of the kernel.
+ Values include:
+
+ - "OpenCL C"
+ - "OpenCL C++"
+ - "HCC"
+ - "OpenMP"
+
+ "LanguageVersion" sequence of - The first integer is the major
+ 2 integers version.
+ - The second integer is the
+ minor version.
+ "Attrs" mapping Mapping of kernel attributes.
+ See
+ :ref:`amdgpu-amdhsa-code-object-kernel-attribute-metadata-mapping-table`
+ for the mapping definition.
+ "Arguments" sequence of Sequence of mappings of the
+ mapping kernel arguments. See
+ :ref:`amdgpu-amdhsa-code-object-kernel-argument-metadata-mapping-table`
+ for the definition of the mapping.
+ "CodeProps" mapping Mapping of properties related to
+ the kernel code. See
+ :ref:`amdgpu-amdhsa-code-object-kernel-code-properties-metadata-mapping-table`
+ for the mapping definition.
+ "DebugProps" mapping Mapping of properties related to
+ the kernel debugging. See
+ :ref:`amdgpu-amdhsa-code-object-kernel-debug-properties-metadata-mapping-table`
+ for the mapping definition.
+ ================= ============== ========= ================================
+
+..
+
+ .. table:: AMDHSA Code Object Kernel Attribute Metadata Mapping
+ :name: amdgpu-amdhsa-code-object-kernel-attribute-metadata-mapping-table
+
+ =================== ============== ========= ==============================
+ String Key Value Type Required? Description
+ =================== ============== ========= ==============================
+ "ReqdWorkGroupSize" sequence of The dispatch work-group size
+ 3 integers X, Y, Z must correspond to the
+ specified values.
+
+ Corresponds to the OpenCL
+ ``reqd_work_group_size``
+ attribute.
+ "WorkGroupSizeHint" sequence of The dispatch work-group size
+ 3 integers X, Y, Z is likely to be the
+ specified values.
+
+ Corresponds to the OpenCL
+ ``work_group_size_hint``
+ attribute.
+ "VecTypeHint" string The name of a scalar or vector
+ type.
+
+ Corresponds to the OpenCL
+ ``vec_type_hint`` attribute.
+ =================== ============== ========= ==============================
+
+..
+
+ .. table:: AMDHSA Code Object Kernel Argument Metadata Mapping
+ :name: amdgpu-amdhsa-code-object-kernel-argument-metadata-mapping-table
+
+ ================= ============== ========= ================================
+ String Key Value Type Required? Description
+ ================= ============== ========= ================================
+ "Name" string Kernel argument name.
+ "TypeName" string Kernel argument type name.
+ "Size" integer Required Kernel argument size in bytes.
+ "Align" integer Required Kernel argument alignment in
+ bytes. Must be a power of two.
+ "ValueKind" string Required Kernel argument kind that
+ specifies how to set up the
+ corresponding argument.
+ Values include:
+
+ "ByValue"
+ The argument is copied
+ directly into the kernarg.
+
+ "GlobalBuffer"
+ A global address space pointer
+ to the buffer data is passed
+ in the kernarg.
+
+ "DynamicSharedPointer"
+ A group address space pointer
+ to dynamically allocated LDS
+ is passed in the kernarg.
+
+ "Sampler"
+ A global address space
+ pointer to a S# is passed in
+ the kernarg.
+
+ "Image"
+ A global address space
+ pointer to a T# is passed in
+ the kernarg.
+
+ "Pipe"
+ A global address space pointer
+ to an OpenCL pipe is passed in
+ the kernarg.
+
+ "Queue"
+ A global address space pointer
+ to an OpenCL device enqueue
+ queue is passed in the
+ kernarg.
+
+ "HiddenGlobalOffsetX"
+ The OpenCL grid dispatch
+ global offset for the X
+ dimension is passed in the
+ kernarg.
+
+ "HiddenGlobalOffsetY"
+ The OpenCL grid dispatch
+ global offset for the Y
+ dimension is passed in the
+ kernarg.
+
+ "HiddenGlobalOffsetZ"
+ The OpenCL grid dispatch
+ global offset for the Z
+ dimension is passed in the
+ kernarg.
+
+ "HiddenNone"
+ An argument that is not used
+ by the kernel. Space needs to
+ be left for it, but it does
+ not need to be set up.
+
+ "HiddenPrintfBuffer"
+ A global address space pointer
+ to the runtime printf buffer
+ is passed in kernarg.
+
+ "HiddenDefaultQueue"
+ A global address space pointer
+ to the OpenCL device enqueue
+ queue that should be used by
+ the kernel by default is
+ passed in the kernarg.
+
+ "HiddenCompletionAction"
+ *TBD*
+
+ .. TODO
+ Add description.
+
+ "ValueType" string Required Kernel argument value type. Only
+ present if "ValueKind" is
+ "ByValue". For vector data
+ types, the value is for the
+ element type. Values include:
+
+ - "Struct"
+ - "I8"
+ - "U8"
+ - "I16"
+ - "U16"
+ - "F16"
+ - "I32"
+ - "U32"
+ - "F32"
+ - "I64"
+ - "U64"
+ - "F64"
+
+ .. TODO
+ How can it be determined if a
+ vector type, and what size
+ vector?
+ "PointeeAlign" integer Alignment in bytes of pointee
+ type for pointer type kernel
+ argument. Must be a power
+ of 2. Only present if
+ "ValueKind" is
+ "DynamicSharedPointer".
+ "AddrSpaceQual" string Kernel argument address space
+ qualifier. Only present if
+ "ValueKind" is "GlobalBuffer" or
+ "DynamicSharedPointer". Values
+ are:
+
+ - "Private"
+ - "Global"
+ - "Constant"
+ - "Local"
+ - "Generic"
+ - "Region"
+
+ .. TODO
+ Is GlobalBuffer only Global
+ or Constant? Is
+ DynamicSharedPointer always
+ Local? Can HCC allow Generic?
+ How can Private or Region
+ ever happen?
+ "AccQual" string Kernel argument access
+ qualifier. Only present if
+ "ValueKind" is "Image" or
+ "Pipe". Values
+ are:
+
+ - "ReadOnly"
+ - "WriteOnly"
+ - "ReadWrite"
+
+ .. TODO
+ Does this apply to
+ GlobalBuffer?
+ "ActualAcc" string The actual memory accesses
+ performed by the kernel on the
+ kernel argument. Only present if
+ "ValueKind" is "GlobalBuffer",
+ "Image", or "Pipe". This may be
+ more restrictive than indicated
+ by "AccQual" to reflect what the
+ kernel actual does. If not
+ present then the runtime must
+ assume what is implied by
+ "AccQual" and "IsConst". Values
+ are:
+
+ - "ReadOnly"
+ - "WriteOnly"
+ - "ReadWrite"
+
+ "IsConst" boolean Indicates if the kernel argument
+ is const qualified. Only present
+ if "ValueKind" is
+ "GlobalBuffer".
+
+ "IsRestrict" boolean Indicates if the kernel argument
+ is restrict qualified. Only
+ present if "ValueKind" is
+ "GlobalBuffer".
+
+ "IsVolatile" boolean Indicates if the kernel argument
+ is volatile qualified. Only
+ present if "ValueKind" is
+ "GlobalBuffer".
+
+ "IsPipe" boolean Indicates if the kernel argument
+ is pipe qualified. Only present
+ if "ValueKind" is "Pipe".
+
+ .. TODO
+ Can GlobalBuffer be pipe
+ qualified?
+ ================= ============== ========= ================================
+
+..
+
+ .. table:: AMDHSA Code Object Kernel Code Properties Metadata Mapping
+ :name: amdgpu-amdhsa-code-object-kernel-code-properties-metadata-mapping-table
+
+ ============================ ============== ========= =====================
+ String Key Value Type Required? Description
+ ============================ ============== ========= =====================
+ "KernargSegmentSize" integer Required The size in bytes of
+ the kernarg segment
+ that holds the values
+ of the arguments to
+ the kernel.
+ "GroupSegmentFixedSize" integer Required The amount of group
+ segment memory
+ required by a
+ work-group in
+ bytes. This does not
+ include any
+ dynamically allocated
+ group segment memory
+ that may be added
+ when the kernel is
+ dispatched.
+ "PrivateSegmentFixedSize" integer Required The amount of fixed
+ private address space
+ memory required for a
+ work-item in
+ bytes. If
+ IsDynamicCallstack
+ is 1 then additional
+ space must be added
+ to this value for the
+ call stack.
+ "KernargSegmentAlign" integer Required The maximum byte
+ alignment of
+ arguments in the
+ kernarg segment. Must
+ be a power of 2.
+ "WavefrontSize" integer Required Wavefront size. Must
+ be a power of 2.
+ "NumSGPRs" integer Number of scalar
+ registers used by a
+ wavefront for
+ GFX6-GFX9. This
+ includes the special
+ SGPRs for VCC, Flat
+ Scratch (GFX7-GFX9)
+ and XNACK (for
+ GFX8-GFX9). It does
+ not include the 16
+ SGPR added if a trap
+ handler is
+ enabled. It is not
+ rounded up to the
+ allocation
+ granularity.
+ "NumVGPRs" integer Number of vector
+ registers used by
+ each work-item for
+ GFX6-GFX9
+ "MaxFlatWorkgroupSize" integer Maximum flat
+ work-group size
+ supported by the
+ kernel in work-items.
+ "IsDynamicCallStack" boolean Indicates if the
+ generated machine
+ code is using a
+ dynamically sized
+ call stack.
+ "IsXNACKEnabled" boolean Indicates if the
+ generated machine
+ code is capable of
+ supporting XNACK.
+ ============================ ============== ========= =====================
+
+..
+
+ .. table:: AMDHSA Code Object Kernel Debug Properties Metadata Mapping
+ :name: amdgpu-amdhsa-code-object-kernel-debug-properties-metadata-mapping-table
+
+ =================================== ============== ========= ==============
+ String Key Value Type Required? Description
+ =================================== ============== ========= ==============
+ "DebuggerABIVersion" string
+ "ReservedNumVGPRs" integer
+ "ReservedFirstVGPR" integer
+ "PrivateSegmentBufferSGPR" integer
+ "WavefrontPrivateSegmentOffsetSGPR" integer
+ =================================== ============== ========= ==============
+
+.. TODO
+ Plan to remove the debug properties metadata.
+
+.. _amdgpu-symbols:
+
+Symbols
+-------
+
+Symbols include the following:
+
+ .. table:: AMDGPU ELF Symbols
+ :name: amdgpu-elf-symbols-table
+
+ ===================== ============== ============= ==================
+ Name Type Section Description
+ ===================== ============== ============= ==================
+ *link-name* ``STT_OBJECT`` - ``.data`` Global variable
+ - ``.rodata``
+ - ``.bss``
+ *link-name*\ ``@kd`` ``STT_OBJECT`` - ``.rodata`` Kernel descriptor
+ *link-name* ``STT_FUNC`` - ``.text`` Kernel entry point
+ ===================== ============== ============= ==================
+
+Global variable
+ Global variables both used and defined by the compilation unit.
+
+ If the symbol is defined in the compilation unit then it is allocated in the
+ appropriate section according to if it has initialized data or is readonly.
+
+ If the symbol is external then its section is ``STN_UNDEF`` and the loader
+ will resolve relocations using the defintion provided by another code object
+ or explicitly defined by the runtime.
+
+ All global symbols, whether defined in the compilation unit or external, are
+ accessed by the machine code indirectly throught a GOT table entry. This
+ allows them to be preemptable. The GOT table is only supported when the target
+ triple OS is ``amdhsa`` (see :ref:`amdgpu-target-triples`).
+
+ .. TODO
+ Add description of linked shared object symbols. Seems undefined symbols
+ are marked as STT_NOTYPE.
+
+Kernel descriptor
+ Every HSA kernel has an associated kernel descriptor. It is the address of the
+ kernel descriptor that is used in the AQL dispatch packet used to invoke the
+ kernel, not the kernel entry point. The layout of the HSA kernel descriptor is
+ defined in :ref:`amdgpu-amdhsa-kernel-descriptor`.
+
+Kernel entry point
+ Every HSA kernel also has a symbol for its machine code entry point.
+
+.. _amdgpu-relocation-records:
+
+Relocation Records
+------------------
+
+AMDGPU backend generates ``Elf64_Rela`` relocation records. Supported
+relocatable fields are:
+
+``word32``
+ This specifies a 32-bit field occupying 4 bytes with arbitrary byte
+ alignment. These values use the same byte order as other word values in the
+ AMD GPU architecture.
+
+``word64``
+ This specifies a 64-bit field occupying 8 bytes with arbitrary byte
+ alignment. These values use the same byte order as other word values in the
+ AMD GPU architecture.
+
+Following notations are used for specifying relocation calculations:
+
+**A**
+ Represents the addend used to compute the value of the relocatable field.
+
+**G**
+ Represents the offset into the global offset table at which the relocation
+ entry’s symbol will reside during execution.
+
+**GOT**
+ Represents the address of the global offset table.
+
+**P**
+ Represents the place (section offset for ``et_rel`` or address for ``et_dyn``)
+ of the storage unit being relocated (computed using ``r_offset``).
+
+**S**
+ Represents the value of the symbol whose index resides in the relocation
+ entry.
+
+The following relocation types are supported:
+
+ .. table:: AMDGPU ELF Relocation Records
+ :name: amdgpu-elf-relocation-records-table
+
+ ========================== ===== ========== ==============================
+ Relocation Type Value Field Calculation
+ ========================== ===== ========== ==============================
+ ``R_AMDGPU_NONE`` 0 *none* *none*
+ ``R_AMDGPU_ABS32_LO`` 1 ``word32`` (S + A) & 0xFFFFFFFF
+ ``R_AMDGPU_ABS32_HI`` 2 ``word32`` (S + A) >> 32
+ ``R_AMDGPU_ABS64`` 3 ``word64`` S + A
+ ``R_AMDGPU_REL32`` 4 ``word32`` S + A - P
+ ``R_AMDGPU_REL64`` 5 ``word64`` S + A - P
+ ``R_AMDGPU_ABS32`` 6 ``word32`` S + A
+ ``R_AMDGPU_GOTPCREL`` 7 ``word32`` G + GOT + A - P
+ ``R_AMDGPU_GOTPCREL32_LO`` 8 ``word32`` (G + GOT + A - P) & 0xFFFFFFFF
+ ``R_AMDGPU_GOTPCREL32_HI`` 9 ``word32`` (G + GOT + A - P) >> 32
+ ``R_AMDGPU_REL32_LO`` 10 ``word32`` (S + A - P) & 0xFFFFFFFF
+ ``R_AMDGPU_REL32_HI`` 11 ``word32`` (S + A - P) >> 32
+ ========================== ===== ========== ==============================
+
+.. _amdgpu-dwarf:
+
+DWARF
+-----
+
+Standard DWARF [DWARF]_ Version 2 sections can be generated. These contain
+information that maps the code object executable code and data to the source
+language constructs. It can be used by tools such as debuggers and profilers.
+
+Address Space Mapping
+~~~~~~~~~~~~~~~~~~~~~
+
+The following address space mapping is used:
+
+ .. table:: AMDGPU DWARF Address Space Mapping
+ :name: amdgpu-dwarf-address-space-mapping-table
+
+ =================== =================
+ DWARF Address Space Memory Space
+ =================== =================
+ 1 Private (Scratch)
+ 2 Local (group/LDS)
+ *omitted* Global
+ *omitted* Constant
+ *omitted* Generic (Flat)
+ *not supported* Region (GDS)
+ =================== =================
+
+See :ref:`amdgpu-address-spaces` for infomration on the memory space terminology
+used in the table.
+
+An ``address_class`` attribute is generated on pointer type DIEs to specify the
+DWARF address space of the value of the pointer when it is in the *private* or
+*local* address space. Otherwise the attribute is omitted.
+
+An ``XDEREF`` operation is generated in location list expressions for variables
+that are allocated in the *private* and *local* address space. Otherwise no
+``XDREF`` is omitted.
+
+Register Mapping
+~~~~~~~~~~~~~~~~
+
+*This section is WIP.*
+
+.. TODO
+ Define DWARF register enumeration.
+
+ If want to present a wavefront state then should expose vector registers as
+ 64 wide (rather than per work-item view that LLVM uses). Either as seperate
+ registers, or a 64x4 byte single register. In either case use a new LANE op
+ (akin to XDREF) to select the current lane usage in a location
+ expression. This would also allow scalar register spilling to vector register
+ lanes to be expressed (currently no debug information is being generated for
+ spilling). If choose a wide single register approach then use LANE in
+ conjunction with PIECE operation to select the dword part of the register for
+ the current lane. If the separate register approach then use LANE to select
+ the register.
+
+Source Text
+~~~~~~~~~~~
+
+*This section is WIP.*
+
+.. TODO
+ DWARF extension to include runtime generated source text.
+
+.. _amdgpu-code-conventions:
+
+Code Conventions
+================
+
+AMDHSA
+------
+
+This section provides code conventions used when the target triple OS is
+``amdhsa`` (see :ref:`amdgpu-target-triples`).
+
+Kernel Dispatch
+~~~~~~~~~~~~~~~
+
+The HSA architected queuing language (AQL) defines a user space memory interface
+that can be used to control the dispatch of kernels, in an agent independent
+way. An agent can have zero or more AQL queues created for it using the ROCm
+runtime, in which AQL packets (all of which are 64 bytes) can be placed. See the
+*HSA Platform System Architecture Specification* [HSA]_ for the AQL queue
+mechanics and packet layouts.
+
+The packet processor of a kernel agent is responsible for detecting and
+dispatching HSA kernels from the AQL queues associated with it. For AMD GPUs the
+packet processor is implemented by the hardware command processor (CP),
+asynchronous dispatch controller (ADC) and shader processor input controller
+(SPI).
+
+The ROCm runtime can be used to allocate an AQL queue object. It uses the kernel
+mode driver to initialize and register the AQL queue with CP.
+
+To dispatch a kernel the following actions are performed. This can occur in the
+CPU host program, or from an HSA kernel executing on a GPU.
+
+1. A pointer to an AQL queue for the kernel agent on which the kernel is to be
+ executed is obtained.
+2. A pointer to the kernel descriptor (see
+ :ref:`amdgpu-amdhsa-kernel-descriptor`) of the kernel to execute is
+ obtained. It must be for a kernel that is contained in a code object that that
+ was loaded by the ROCm runtime on the kernel agent with which the AQL queue is
+ associated.
+3. Space is allocated for the kernel arguments using the ROCm runtime allocator
+ for a memory region with the kernarg property for the kernel agent that will
+ execute the kernel. It must be at least 16 byte aligned.
+4. Kernel argument values are assigned to the kernel argument memory
+ allocation. The layout is defined in the *HSA Programmer’s Language Reference*
+ [HSA]_. For AMDGPU the kernel execution directly accesses the kernel argument
+ memory in the same way constant memory is accessed. (Note that the HSA
+ specification allows an implementation to copy the kernel argument contents to
+ another location that is accessed by the kernel.)
+5. An AQL kernel dispatch packet is created on the AQL queue. The ROCm runtime
+ api uses 64 bit atomic operations to reserve space in the AQL queue for the
+ packet. The packet must be set up, and the final write must use an atomic
+ store release to set the packet kind to ensure the packet contents are
+ visible to the kernel agent. AQL defines a doorbell signal mechanism to
+ notify the kernel agent that the AQL queue has been updated. These rules, and
+ the layout of the AQL queue and kernel dispatch packet is defined in the *HSA
+ System Architecture Specification* [HSA]_.
+6. A kernel dispatch packet includes information about the actual dispatch,
+ such as grid and work-group size, together with information from the code
+ object about the kernel, such as segment sizes. The ROCm runtime queries on
+ the kernel symbol can be used to obtain the code object values which are
+ recorded in the :ref:`amdgpu-code-object-metadata`.
+7. CP executes micro-code and is responsible for detecting and setting up the
+ GPU to execute the wavefronts of a kernel dispatch.
+8. CP ensures that when the a wavefront starts executing the kernel machine
+ code, the scalar general purpose registers (SGPR) and vector general purpose
+ registers (VGPR) are set up as required by the machine code. The required
+ setup is defined in the :ref:`amdgpu-amdhsa-kernel-descriptor`. The initial
+ register state is defined in
+ :ref:`amdgpu-amdhsa-initial-kernel-execution-state`.
+9. The prolog of the kernel machine code (see
+ :ref:`amdgpu-amdhsa-kernel-prolog`) sets up the machine state as necessary
+ before continuing executing the machine code that corresponds to the kernel.
+10. When the kernel dispatch has completed execution, CP signals the completion
+ signal specified in the kernel dispatch packet if not 0.
+
+.. _amdgpu-amdhsa-memory-spaces:
+
+Memory Spaces
+~~~~~~~~~~~~~
+
+The memory space properties are:
+
+ .. table:: AMDHSA Memory Spaces
+ :name: amdgpu-amdhsa-memory-spaces-table
+
+ ================= =========== ======== ======= ==================
+ Memory Space Name HSA Segment Hardware Address NULL Value
+ Name Name Size
+ ================= =========== ======== ======= ==================
+ Private private scratch 32 0x00000000
+ Local group LDS 32 0xFFFFFFFF
+ Global global global 64 0x0000000000000000
+ Constant constant *same as 64 0x0000000000000000
+ global*
+ Generic flat flat 64 0x0000000000000000
+ Region N/A GDS 32 *not implemented
+ for AMDHSA*
+ ================= =========== ======== ======= ==================
+
+The global and constant memory spaces both use global virtual addresses, which
+are the same virtual address space used by the CPU. However, some virtual
+addresses may only be accessible to the CPU, some only accessible by the GPU,
+and some by both.
+
+Using the constant memory space indicates that the data will not change during
+the execution of the kernel. This allows scalar read instructions to be
+used. The vector and scalar L1 caches are invalidated of volatile data before
+each kernel dispatch execution to allow constant memory to change values between
+kernel dispatches.
+
+The local memory space uses the hardware Local Data Store (LDS) which is
+automatically allocated when the hardware creates work-groups of wavefronts, and
+freed when all the wavefronts of a work-group have terminated. The data store
+(DS) instructions can be used to access it.
+
+The private memory space uses the hardware scratch memory support. If the kernel
+uses scratch, then the hardware allocates memory that is accessed using
+wavefront lane dword (4 byte) interleaving. The mapping used from private
+address to physical address is:
+
+ ``wavefront-scratch-base +
+ (private-address * wavefront-size * 4) +
+ (wavefront-lane-id * 4)``
+
+There are different ways that the wavefront scratch base address is determined
+by a wavefront (see :ref:`amdgpu-amdhsa-initial-kernel-execution-state`). This
+memory can be accessed in an interleaved manner using buffer instruction with
+the scratch buffer descriptor and per wave scratch offset, by the scratch
+instructions, or by flat instructions. If each lane of a wavefront accesses the
+same private address, the interleaving results in adjacent dwords being accessed
+and hence requires fewer cache lines to be fetched. Multi-dword access is not
+supported except by flat and scratch instructions in GFX9.
+
+The generic address space uses the hardware flat address support available in
+GFX7-GFX9. This uses two fixed ranges of virtual addresses (the private and
+local appertures), that are outside the range of addressible global memory, to
+map from a flat address to a private or local address.
+
+FLAT instructions can take a flat address and access global, private (scratch)
+and group (LDS) memory depending in if the address is within one of the
+apperture ranges. Flat access to scratch requires hardware aperture setup and
+setup in the kernel prologue (see :ref:`amdgpu-amdhsa-flat-scratch`). Flat
+access to LDS requires hardware aperture setup and M0 (GFX7-GFX8) register setup
+(see :ref:`amdgpu-amdhsa-m0`).
+
+To convert between a segment address and a flat address the base address of the
+appertures address can be used. For GFX7-GFX8 these are available in the
+:ref:`amdgpu-amdhsa-hsa-aql-queue` the address of which can be obtained with
+Queue Ptr SGPR (see :ref:`amdgpu-amdhsa-initial-kernel-execution-state`). For
+GFX9 the appature base addresses are directly available as inline constant
+registers ``SRC_SHARED_BASE/LIMIT`` and ``SRC_PRIVATE_BASE/LIMIT``. In 64 bit
+address mode the apperture sizes are 2^32 bytes and the base is aligned to 2^32
+which makes it easier to convert from flat to segment or segment to flat.
+
+HSA Image and Samplers
+~~~~~~~~~~~~~~~~~~~~~~
+
+Image and sample handles created by the ROCm runtime are 64 bit addresses of a
+hardware 32 byte V# and 48 byte S# object respectively. In order to support the
+HSA ``query_sampler`` operations two extra dwords are used to store the HSA BRIG
+enumeration values for the queries that are not trivially deducible from the S#
+representation.
+
+HSA Signals
+~~~~~~~~~~~
+
+Signal handles created by the ROCm runtime are 64 bit addresses of a structure
+allocated in memory accessible from both the CPU and GPU. The structure is
+defined by the ROCm runtime and subject to change between releases (see
+[AMD-ROCm-github]_).
+
+.. _amdgpu-amdhsa-hsa-aql-queue:
+
+HSA AQL Queue
+~~~~~~~~~~~~~
+
+The AQL queue structure is defined by the ROCm runtime and subject to change
+between releases (see [AMD-ROCm-github]_). For some processors it contains
+fields needed to implement certain language features such as the flat address
+aperture bases. It also contains fields used by CP such as managing the
+allocation of scratch memory.
+
+.. _amdgpu-amdhsa-kernel-descriptor:
+
+Kernel Descriptor
+~~~~~~~~~~~~~~~~~
+
+A kernel descriptor consists of the information needed by CP to initiate the
+execution of a kernel, including the entry point address of the machine code
+that implements the kernel.
+
+Kernel Descriptor for GFX6-GFX9
++++++++++++++++++++++++++++++++
+
+CP microcode requires the Kernel descritor to be allocated on 64 byte alignment.
+
+ .. table:: Kernel Descriptor for GFX6-GFX9
+ :name: amdgpu-amdhsa-kernel-descriptor-gfx6-gfx9-table
+
+ ======= ======= =============================== ===========================
+ Bits Size Field Name Description
+ ======= ======= =============================== ===========================
+ 31:0 4 bytes group_segment_fixed_size The amount of fixed local
+ address space memory
+ required for a work-group
+ in bytes. This does not
+ include any dynamically
+ allocated local address
+ space memory that may be
+ added when the kernel is
+ dispatched.
+ 63:32 4 bytes private_segment_fixed_size The amount of fixed
+ private address space
+ memory required for a
+ work-item in bytes. If
+ is_dynamic_callstack is 1
+ then additional space must
+ be added to this value for
+ the call stack.
+ 95:64 4 bytes max_flat_workgroup_size Maximum flat work-group
+ size supported by the
+ kernel in work-items.
+ 96 1 bit is_dynamic_call_stack Indicates if the generated
+ machine code is using a
+ dynamically sized call
+ stack.
+ 97 1 bit is_xnack_enabled Indicates if the generated
+ machine code is capable of
+ suppoting XNACK.
+ 127:98 30 bits Reserved. Must be 0.
+ 191:128 8 bytes kernel_code_entry_byte_offset Byte offset (possibly
+ negative) from base
+ address of kernel
+ descriptor to kernel's
+ entry point instruction
+ which must be 256 byte
+ aligned.
+ 383:192 24 Reserved. Must be 0.
+ bytes
+ 415:384 4 bytes compute_pgm_rsrc1 Compute Shader (CS)
+ program settings used by
+ CP to set up
+ ``COMPUTE_PGM_RSRC1``
+ configuration
+ register. See
+ :ref:`amdgpu-amdhsa-compute_pgm_rsrc1_t-gfx6-gfx9-table`.
+ 447:416 4 bytes compute_pgm_rsrc2 Compute Shader (CS)
+ program settings used by
+ CP to set up
+ ``COMPUTE_PGM_RSRC2``
+ configuration
+ register. See
+ :ref:`amdgpu-amdhsa-compute_pgm_rsrc2-gfx6-gfx9-table`.
+ 448 1 bit enable_sgpr_private_segment Enable the setup of the
+ _buffer SGPR user data registers
+ (see
+ :ref:`amdgpu-amdhsa-initial-kernel-execution-state`).
+
+ The total number of SGPR
+ user data registers
+ requested must not exceed
+ 16 and match value in
+ ``compute_pgm_rsrc2.user_sgpr.user_sgpr_count``.
+ Any requests beyond 16
+ will be ignored.
+ 449 1 bit enable_sgpr_dispatch_ptr *see above*
+ 450 1 bit enable_sgpr_queue_ptr *see above*
+ 451 1 bit enable_sgpr_kernarg_segment_ptr *see above*
+ 452 1 bit enable_sgpr_dispatch_id *see above*
+ 453 1 bit enable_sgpr_flat_scratch_init *see above*
+ 454 1 bit enable_sgpr_private_segment *see above*
+ _size
+ 455 1 bit enable_sgpr_grid_workgroup Not implemented in CP and
+ _count_X should always be 0.
+ 456 1 bit enable_sgpr_grid_workgroup Not implemented in CP and
+ _count_Y should always be 0.
+ 457 1 bit enable_sgpr_grid_workgroup Not implemented in CP and
+ _count_Z should always be 0.
+ 463:458 6 bits Reserved. Must be 0.
+ 511:464 4 Reserved. Must be 0.
+ bytes
+ 512 **Total size 64 bytes.**
+ ======= ===================================================================
+
+..
+
+ .. table:: compute_pgm_rsrc1 for GFX6-GFX9
+ :name: amdgpu-amdhsa-compute_pgm_rsrc1_t-gfx6-gfx9-table
+
+ ======= ======= =============================== ===========================================================================
+ Bits Size Field Name Description
+ ======= ======= =============================== ===========================================================================
+ 5:0 6 bits granulated_workitem_vgpr_count Number of vector registers
+ used by each work-item,
+ granularity is device
+ specific:
+
+ GFX6-9
+ roundup((max-vgpg + 1)
+ / 4) - 1
+
+ Used by CP to set up
+ ``COMPUTE_PGM_RSRC1.VGPRS``.
+ 9:6 4 bits granulated_wavefront_sgpr_count Number of scalar registers
+ used by a wavefront,
+ granularity is device
+ specific:
+
+ GFX6-8
+ roundup((max-sgpg + 1)
+ / 8) - 1
+ GFX9
+ roundup((max-sgpg + 1)
+ / 16) - 1
+
+ Includes the special SGPRs
+ for VCC, Flat Scratch (for
+ GFX7 onwards) and XNACK
+ (for GFX8 onwards). It does
+ not include the 16 SGPR
+ added if a trap handler is
+ enabled.
+
+ Used by CP to set up
+ ``COMPUTE_PGM_RSRC1.SGPRS``.
+ 11:10 2 bits priority Must be 0.
+
+ Start executing wavefront
+ at the specified priority.
+
+ CP is responsible for
+ filling in
+ ``COMPUTE_PGM_RSRC1.PRIORITY``.
+ 13:12 2 bits float_mode_round_32 Wavefront starts execution
+ with specified rounding
+ mode for single (32
+ bit) floating point
+ precision floating point
+ operations.
+
+ Floating point rounding
+ mode values are defined in
+ :ref:`amdgpu-amdhsa-floating-point-rounding-mode-enumeration-values-table`.
+
+ Used by CP to set up
+ ``COMPUTE_PGM_RSRC1.FLOAT_MODE``.
+ 15:14 2 bits float_mode_round_16_64 Wavefront starts execution
+ with specified rounding
+ denorm mode for half/double (16
+ and 64 bit) floating point
+ precision floating point
+ operations.
+
+ Floating point rounding
+ mode values are defined in
+ :ref:`amdgpu-amdhsa-floating-point-rounding-mode-enumeration-values-table`.
+
+ Used by CP to set up
+ ``COMPUTE_PGM_RSRC1.FLOAT_MODE``.
+ 17:16 2 bits float_mode_denorm_32 Wavefront starts execution
+ with specified denorm mode
+ for single (32
+ bit) floating point
+ precision floating point
+ operations.
+
+ Floating point denorm mode
+ values are defined in
+ :ref:`amdgpu-amdhsa-floating-point-denorm-mode-enumeration-values-table`.
+
+ Used by CP to set up
+ ``COMPUTE_PGM_RSRC1.FLOAT_MODE``.
+ 19:18 2 bits float_mode_denorm_16_64 Wavefront starts execution
+ with specified denorm mode
+ for half/double (16
+ and 64 bit) floating point
+ precision floating point
+ operations.
+
+ Floating point denorm mode
+ values are defined in
+ :ref:`amdgpu-amdhsa-floating-point-denorm-mode-enumeration-values-table`.
+
+ Used by CP to set up
+ ``COMPUTE_PGM_RSRC1.FLOAT_MODE``.
+ 20 1 bit priv Must be 0.
+
+ Start executing wavefront
+ in privilege trap handler
+ mode.
+
+ CP is responsible for
+ filling in
+ ``COMPUTE_PGM_RSRC1.PRIV``.
+ 21 1 bit enable_dx10_clamp Wavefront starts execution
+ with DX10 clamp mode
+ enabled. Used by the vector
+ ALU to force DX-10 style
+ treatment of NaN's (when
+ set, clamp NaN to zero,
+ otherwise pass NaN
+ through).
+
+ Used by CP to set up
+ ``COMPUTE_PGM_RSRC1.DX10_CLAMP``.
+ 22 1 bit debug_mode Must be 0.
+
+ Start executing wavefront
+ in single step mode.
+
+ CP is responsible for
+ filling in
+ ``COMPUTE_PGM_RSRC1.DEBUG_MODE``.
+ 23 1 bit enable_ieee_mode Wavefront starts execution
+ with IEEE mode
+ enabled. Floating point
+ opcodes that support
+ exception flag gathering
+ will quiet and propagate
+ signaling-NaN inputs per
+ IEEE 754-2008. Min_dx10 and
+ max_dx10 become IEEE
+ 754-2008 compliant due to
+ signaling-NaN propagation
+ and quieting.
+
+ Used by CP to set up
+ ``COMPUTE_PGM_RSRC1.IEEE_MODE``.
+ 24 1 bit bulky Must be 0.
+
+ Only one work-group allowed
+ to execute on a compute
+ unit.
+
+ CP is responsible for
+ filling in
+ ``COMPUTE_PGM_RSRC1.BULKY``.
+ 25 1 bit cdbg_user Must be 0.
+
+ Flag that can be used to
+ control debugging code.
+
+ CP is responsible for
+ filling in
+ ``COMPUTE_PGM_RSRC1.CDBG_USER``.
+ 31:26 6 bits Reserved. Must be 0.
+ 32 **Total size 4 bytes**
+ ======= ===================================================================================================================
+
+..
+
+ .. table:: compute_pgm_rsrc2 for GFX6-GFX9
+ :name: amdgpu-amdhsa-compute_pgm_rsrc2-gfx6-gfx9-table
+
+ ======= ======= =============================== ===========================================================================
+ Bits Size Field Name Description
+ ======= ======= =============================== ===========================================================================
+ 0 1 bit enable_sgpr_private_segment Enable the setup of the
+ _wave_offset SGPR wave scratch offset
+ system register (see
+ :ref:`amdgpu-amdhsa-initial-kernel-execution-state`).
+
+ Used by CP to set up
+ ``COMPUTE_PGM_RSRC2.SCRATCH_EN``.
+ 5:1 5 bits user_sgpr_count The total number of SGPR
+ user data registers
+ requested. This number must
+ match the number of user
+ data registers enabled.
+
+ Used by CP to set up
+ ``COMPUTE_PGM_RSRC2.USER_SGPR``.
+ 6 1 bit enable_trap_handler Set to 1 if code contains a
+ TRAP instruction which
+ requires a trap hander to
+ be enabled.
+
+ CP sets
+ ``COMPUTE_PGM_RSRC2.TRAP_PRESENT``
+ if the runtime has
+ installed a trap handler
+ regardless of the setting
+ of this field.
+ 7 1 bit enable_sgpr_workgroup_id_x Enable the setup of the
+ system SGPR register for
+ the work-group id in the X
+ dimension (see
+ :ref:`amdgpu-amdhsa-initial-kernel-execution-state`).
+
+ Used by CP to set up
+ ``COMPUTE_PGM_RSRC2.TGID_X_EN``.
+ 8 1 bit enable_sgpr_workgroup_id_y Enable the setup of the
+ system SGPR register for
+ the work-group id in the Y
+ dimension (see
+ :ref:`amdgpu-amdhsa-initial-kernel-execution-state`).
+
+ Used by CP to set up
+ ``COMPUTE_PGM_RSRC2.TGID_Y_EN``.
+ 9 1 bit enable_sgpr_workgroup_id_z Enable the setup of the
+ system SGPR register for
+ the work-group id in the Z
+ dimension (see
+ :ref:`amdgpu-amdhsa-initial-kernel-execution-state`).
+
+ Used by CP to set up
+ ``COMPUTE_PGM_RSRC2.TGID_Z_EN``.
+ 10 1 bit enable_sgpr_workgroup_info Enable the setup of the
+ system SGPR register for
+ work-group information (see
+ :ref:`amdgpu-amdhsa-initial-kernel-execution-state`).
+
+ Used by CP to set up
+ ``COMPUTE_PGM_RSRC2.TGID_SIZE_EN``.
+ 12:11 2 bits enable_vgpr_workitem_id Enable the setup of the
+ VGPR system registers used
+ for the work-item ID.
+ :ref:`amdgpu-amdhsa-system-vgpr-work-item-id-enumeration-values-table`
+ defines the values.
+
+ Used by CP to set up
+ ``COMPUTE_PGM_RSRC2.TIDIG_CMP_CNT``.
+ 13 1 bit enable_exception_address_watch Must be 0.
+
+ Wavefront starts execution
+ with address watch
+ exceptions enabled which
+ are generated when L1 has
+ witnessed a thread access
+ an *address of
+ interest*.
+
+ CP is responsible for
+ filling in the address
+ watch bit in
+ ``COMPUTE_PGM_RSRC2.EXCP_EN_MSB``
+ according to what the
+ runtime requests.
+ 14 1 bit enable_exception_memory Must be 0.
+
+ Wavefront starts execution
+ with memory violation
+ exceptions exceptions
+ enabled which are generated
+ when a memory violation has
+ occurred for this wave from
+ L1 or LDS
+ (write-to-read-only-memory,
+ mis-aligned atomic, LDS
+ address out of range,
+ illegal address, etc.).
+
+ CP sets the memory
+ violation bit in
+ ``COMPUTE_PGM_RSRC2.EXCP_EN_MSB``
+ according to what the
+ runtime requests.
+ 23:15 9 bits granulated_lds_size Must be 0.
+
+ CP uses the rounded value
+ from the dispatch packet,
+ not this value, as the
+ dispatch may contain
+ dynamically allocated group
+ segment memory. CP writes
+ directly to
+ ``COMPUTE_PGM_RSRC2.LDS_SIZE``.
+
+ Amount of group segment
+ (LDS) to allocate for each
+ work-group. Granularity is
+ device specific:
+
+ GFX6:
+ roundup(lds-size / (64 * 4))
+ GFX7-GFX9:
+ roundup(lds-size / (128 * 4))
+
+ 24 1 bit enable_exception_ieee_754_fp Wavefront starts execution
+ _invalid_operation with specified exceptions
+ enabled.
+
+ Used by CP to set up
+ ``COMPUTE_PGM_RSRC2.EXCP_EN``
+ (set from bits 0..6).
+
+ IEEE 754 FP Invalid
+ Operation
+ 25 1 bit enable_exception_fp_denormal FP Denormal one or more
+ _source input operands is a
+ denormal number
+ 26 1 bit enable_exception_ieee_754_fp IEEE 754 FP Division by
+ _division_by_zero Zero
+ 27 1 bit enable_exception_ieee_754_fp IEEE 754 FP FP Overflow
+ _overflow
+ 28 1 bit enable_exception_ieee_754_fp IEEE 754 FP Underflow
+ _underflow
+ 29 1 bit enable_exception_ieee_754_fp IEEE 754 FP Inexact
+ _inexact
+ 30 1 bit enable_exception_int_divide_by Integer Division by Zero
+ _zero (rcp_iflag_f32 instruction
+ only)
+ 31 1 bit Reserved. Must be 0.
+ 32 **Total size 4 bytes.**
+ ======= ===================================================================================================================
+
+..
+
+ .. table:: Floating Point Rounding Mode Enumeration Values
+ :name: amdgpu-amdhsa-floating-point-rounding-mode-enumeration-values-table
+
+ ===================================== ===== ===============================
+ Enumeration Name Value Description
+ ===================================== ===== ===============================
+ AMD_FLOAT_ROUND_MODE_NEAR_EVEN 0 Round Ties To Even
+ AMD_FLOAT_ROUND_MODE_PLUS_INFINITY 1 Round Toward +infinity
+ AMD_FLOAT_ROUND_MODE_MINUS_INFINITY 2 Round Toward -infinity
+ AMD_FLOAT_ROUND_MODE_ZERO 3 Round Toward 0
+ ===================================== ===== ===============================
+
+..
+
+ .. table:: Floating Point Denorm Mode Enumeration Values
+ :name: amdgpu-amdhsa-floating-point-denorm-mode-enumeration-values-table
+
+ ===================================== ===== ===============================
+ Enumeration Name Value Description
+ ===================================== ===== ===============================
+ AMD_FLOAT_DENORM_MODE_FLUSH_SRC_DST 0 Flush Source and Destination
+ Denorms
+ AMD_FLOAT_DENORM_MODE_FLUSH_DST 1 Flush Output Denorms
+ AMD_FLOAT_DENORM_MODE_FLUSH_SRC 2 Flush Source Denorms
+ AMD_FLOAT_DENORM_MODE_FLUSH_NONE 3 No Flush
+ ===================================== ===== ===============================
+
+..
+
+ .. table:: System VGPR Work-Item ID Enumeration Values
+ :name: amdgpu-amdhsa-system-vgpr-work-item-id-enumeration-values-table
+
+ ===================================== ===== ===============================
+ Enumeration Name Value Description
+ ===================================== ===== ===============================
+ AMD_SYSTEM_VGPR_WORKITEM_ID_X 0 Set work-item X dimension ID.
+ AMD_SYSTEM_VGPR_WORKITEM_ID_X_Y 1 Set work-item X and Y
+ dimensions ID.
+ AMD_SYSTEM_VGPR_WORKITEM_ID_X_Y_Z 2 Set work-item X, Y and Z
+ dimensions ID.
+ AMD_SYSTEM_VGPR_WORKITEM_ID_UNDEFINED 3 Undefined.
+ ===================================== ===== ===============================
+
+.. _amdgpu-amdhsa-initial-kernel-execution-state:
+
+Initial Kernel Execution State
+~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
+
+This section defines the register state that will be set up by the packet
+processor prior to the start of execution of every wavefront. This is limited by
+the constraints of the hardware controllers of CP/ADC/SPI.
+
+The order of the SGPR registers is defined, but the compiler can specify which
+ones are actually setup in the kernel descriptor using the ``enable_sgpr_*`` bit
+fields (see :ref:`amdgpu-amdhsa-kernel-descriptor`). The register numbers used
+for enabled registers are dense starting at SGPR0: the first enabled register is
+SGPR0, the next enabled register is SGPR1 etc.; disabled registers do not have
+an SGPR number.
+
+The initial SGPRs comprise up to 16 User SRGPs that are set by CP and apply to
+all waves of the grid. It is possible to specify more than 16 User SGPRs using
+the ``enable_sgpr_*`` bit fields, in which case only the first 16 are actually
+initialized. These are then immediately followed by the System SGPRs that are
+set up by ADC/SPI and can have different values for each wave of the grid
+dispatch.
+
+SGPR register initial state is defined in
+:ref:`amdgpu-amdhsa-sgpr-register-set-up-order-table`.
+
+ .. table:: SGPR Register Set Up Order
+ :name: amdgpu-amdhsa-sgpr-register-set-up-order-table
+
+ ========== ========================== ====== ==============================
+ SGPR Order Name Number Description
+ (kernel descriptor enable of
+ field) SGPRs
+ ========== ========================== ====== ==============================
+ First Private Segment Buffer 4 V# that can be used, together
+ (enable_sgpr_private with Scratch Wave Offset as an
+ _segment_buffer) offset, to access the private
+ memory space using a segment
+ address.
+
+ CP uses the value provided by
+ the runtime.
+ then Dispatch Ptr 2 64 bit address of AQL dispatch
+ (enable_sgpr_dispatch_ptr) packet for kernel dispatch
+ actually executing.
+ then Queue Ptr 2 64 bit address of amd_queue_t
+ (enable_sgpr_queue_ptr) object for AQL queue on which
+ the dispatch packet was
+ queued.
+ then Kernarg Segment Ptr 2 64 bit address of Kernarg
+ (enable_sgpr_kernarg segment. This is directly
+ _segment_ptr) copied from the
+ kernarg_address in the kernel
+ dispatch packet.
+
+ Having CP load it once avoids
+ loading it at the beginning of
+ every wavefront.
+ then Dispatch Id 2 64 bit Dispatch ID of the
+ (enable_sgpr_dispatch_id) dispatch packet being
+ executed.
+ then Flat Scratch Init 2 This is 2 SGPRs:
+ (enable_sgpr_flat_scratch
+ _init) GFX6
+ Not supported.
+ GFX7-GFX8
+ The first SGPR is a 32 bit
+ byte offset from
+ ``SH_HIDDEN_PRIVATE_BASE_VIMID``
+ to per SPI base of memory
+ for scratch for the queue
+ executing the kernel
+ dispatch. CP obtains this
+ from the runtime.
+
+ This is the same offset used
+ in computing the Scratch
+ Segment Buffer base
+ address. The value of
+ Scratch Wave Offset must be
+ added by the kernel machine
+ code and moved to SGPRn-4
+ for use as the FLAT SCRATCH
+ BASE in flat memory
+ instructions.
+
+ The second SGPR is 32 bit
+ byte size of a single
+ work-item’s scratch memory
+ usage. This is directly
+ loaded from the kernel
+ dispatch packet Private
+ Segment Byte Size and
+ rounded up to a multiple of
+ DWORD.
+
+ The kernel code must move to
+ SGPRn-3 for use as the FLAT
+ SCRATCH SIZE in flat memory
+ instructions. Having CP load
+ it once avoids loading it at
+ the beginning of every
+ wavefront.
+ GFX9
+ This is the 64 bit base
+ address of the per SPI
+ scratch backing memory
+ managed by SPI for the queue
+ executing the kernel
+ dispatch. CP obtains this
+ from the runtime (and
+ divides it if there are
+ multiple Shader Arrays each
+ with its own SPI). The value
+ of Scratch Wave Offset must
+ be added by the kernel
+ machine code and moved to
+ SGPRn-4 and SGPRn-3 for use
+ as the FLAT SCRATCH BASE in
+ flat memory instructions.
+ then Private Segment Size 1 The 32 bit byte size of a
+ (enable_sgpr_private single work-item’s scratch
+ _segment_size) memory allocation. This is the
+ value from the kernel dispatch
+ packet Private Segment Byte
+ Size rounded up by CP to a
+ multiple of DWORD.
+
+ Having CP load it once avoids
+ loading it at the beginning of
+ every wavefront.
+
+ This is not used for
+ GFX7-GFX8 since it is the same
+ value as the second SGPR of
+ Flat Scratch Init. However, it
+ may be needed for GFX9 which
+ changes the meaning of the
+ Flat Scratch Init value.
+ then Grid Work-Group Count X 1 32 bit count of the number of
+ (enable_sgpr_grid work-groups in the X dimension
+ _workgroup_count_X) for the grid being
+ executed. Computed from the
+ fields in the kernel dispatch
+ packet as ((grid_size.x +
+ workgroup_size.x - 1) /
+ workgroup_size.x).
+ then Grid Work-Group Count Y 1 32 bit count of the number of
+ (enable_sgpr_grid work-groups in the Y dimension
+ _workgroup_count_Y && for the grid being
+ less than 16 previous executed. Computed from the
+ SGPRs) fields in the kernel dispatch
+ packet as ((grid_size.y +
+ workgroup_size.y - 1) /
+ workgroupSize.y).
+
+ Only initialized if <16
+ previous SGPRs initialized.
+ then Grid Work-Group Count Z 1 32 bit count of the number of
+ (enable_sgpr_grid work-groups in the Z dimension
+ _workgroup_count_Z && for the grid being
+ less than 16 previous executed. Computed from the
+ SGPRs) fields in the kernel dispatch
+ packet as ((grid_size.z +
+ workgroup_size.z - 1) /
+ workgroupSize.z).
+
+ Only initialized if <16
+ previous SGPRs initialized.
+ then Work-Group Id X 1 32 bit work-group id in X
+ (enable_sgpr_workgroup_id dimension of grid for
+ _X) wavefront.
+ then Work-Group Id Y 1 32 bit work-group id in Y
+ (enable_sgpr_workgroup_id dimension of grid for
+ _Y) wavefront.
+ then Work-Group Id Z 1 32 bit work-group id in Z
+ (enable_sgpr_workgroup_id dimension of grid for
+ _Z) wavefront.
+ then Work-Group Info 1 {first_wave, 14’b0000,
+ (enable_sgpr_workgroup ordered_append_term[10:0],
+ _info) threadgroup_size_in_waves[5:0]}
+ then Scratch Wave Offset 1 32 bit byte offset from base
+ (enable_sgpr_private of scratch base of queue
+ _segment_wave_offset) executing the kernel
+ dispatch. Must be used as an
+ offset with Private
+ segment address when using
+ Scratch Segment Buffer. It
+ must be used to set up FLAT
+ SCRATCH for flat addressing
+ (see
+ :ref:`amdgpu-amdhsa-flat-scratch`).
+ ========== ========================== ====== ==============================
+
+The order of the VGPR registers is defined, but the compiler can specify which
+ones are actually setup in the kernel descriptor using the ``enable_vgpr*`` bit
+fields (see :ref:`amdgpu-amdhsa-kernel-descriptor`). The register numbers used
+for enabled registers are dense starting at VGPR0: the first enabled register is
+VGPR0, the next enabled register is VGPR1 etc.; disabled registers do not have a
+VGPR number.
+
+VGPR register initial state is defined in
+:ref:`amdgpu-amdhsa-vgpr-register-set-up-order-table`.
+
+ .. table:: VGPR Register Set Up Order
+ :name: amdgpu-amdhsa-vgpr-register-set-up-order-table
+
+ ========== ========================== ====== ==============================
+ VGPR Order Name Number Description
+ (kernel descriptor enable of
+ field) VGPRs
+ ========== ========================== ====== ==============================
+ First Work-Item Id X 1 32 bit work item id in X
+ (Always initialized) dimension of work-group for
+ wavefront lane.
+ then Work-Item Id Y 1 32 bit work item id in Y
+ (enable_vgpr_workitem_id dimension of work-group for
+ > 0) wavefront lane.
+ then Work-Item Id Z 1 32 bit work item id in Z
+ (enable_vgpr_workitem_id dimension of work-group for
+ > 1) wavefront lane.
+ ========== ========================== ====== ==============================
+
+The setting of registers is is done by GPU CP/ADC/SPI hardware as follows:
+
+1. SGPRs before the Work-Group Ids are set by CP using the 16 User Data
+ registers.
+2. Work-group Id registers X, Y, Z are set by ADC which supports any
+ combination including none.
+3. Scratch Wave Offset is set by SPI in a per wave basis which is why its value
+ cannot included with the flat scratch init value which is per queue.
+4. The VGPRs are set by SPI which only supports specifying either (X), (X, Y)
+ or (X, Y, Z).
+
+Flat Scratch register pair are adjacent SGRRs so they can be moved as a 64 bit
+value to the hardware required SGPRn-3 and SGPRn-4 respectively.
+
+The global segment can be accessed either using buffer instructions (GFX6 which
+has V# 64 bit address support), flat instructions (GFX7-9), or global
+instructions (GFX9).
+
+If buffer operations are used then the compiler can generate a V# with the
+following properties:
+
+* base address of 0
+* no swizzle
+* ATC: 1 if IOMMU present (such as APU)
+* ptr64: 1
+* MTYPE set to support memory coherence that matches the runtime (such as CC for
+ APU and NC for dGPU).
+
+.. _amdgpu-amdhsa-kernel-prolog:
+
+Kernel Prolog
+~~~~~~~~~~~~~
+
+.. _amdgpu-amdhsa-m0:
+
+M0
+++
+
+GFX6-GFX8
+ The M0 register must be initialized with a value at least the total LDS size
+ if the kernel may access LDS via DS or flat operations. Total LDS size is
+ available in dispatch packet. For M0, it is also possible to use maximum
+ possible value of LDS for given target (0x7FFF for GFX6 and 0xFFFF for
+ GFX7-GFX8).
+GFX9
+ The M0 register is not used for range checking LDS accesses and so does not
+ need to be initialized in the prolog.
+
+.. _amdgpu-amdhsa-flat-scratch:
+
+Flat Scratch
+++++++++++++
+
+If the kernel may use flat operations to access scratch memory, the prolog code
+must set up FLAT_SCRATCH register pair (FLAT_SCRATCH_LO/FLAT_SCRATCH_HI which
+are in SGPRn-4/SGPRn-3). Initialization uses Flat Scratch Init and Scratch Wave
+Offset SGPR registers (see :ref:`amdgpu-amdhsa-initial-kernel-execution-state`):
+
+GFX6
+ Flat scratch is not supported.
+
+GFX7-8
+ 1. The low word of Flat Scratch Init is 32 bit byte offset from
+ ``SH_HIDDEN_PRIVATE_BASE_VIMID`` to the base of scratch backing memory
+ being managed by SPI for the queue executing the kernel dispatch. This is
+ the same value used in the Scratch Segment Buffer V# base address. The
+ prolog must add the value of Scratch Wave Offset to get the wave's byte
+ scratch backing memory offset from ``SH_HIDDEN_PRIVATE_BASE_VIMID``. Since
+ FLAT_SCRATCH_LO is in units of 256 bytes, the offset must be right shifted
+ by 8 before moving into FLAT_SCRATCH_LO.
+ 2. The second word of Flat Scratch Init is 32 bit byte size of a single
+ work-items scratch memory usage. This is directly loaded from the kernel
+ dispatch packet Private Segment Byte Size and rounded up to a multiple of
+ DWORD. Having CP load it once avoids loading it at the beginning of every
+ wavefront. The prolog must move it to FLAT_SCRATCH_LO for use as FLAT SCRATCH
+ SIZE.
+GFX9
+ The Flat Scratch Init is the 64 bit address of the base of scratch backing
+ memory being managed by SPI for the queue executing the kernel dispatch. The
+ prolog must add the value of Scratch Wave Offset and moved to the FLAT_SCRATCH
+ pair for use as the flat scratch base in flat memory instructions.
+
+.. _amdgpu-amdhsa-memory-model:
+
+Memory Model
+~~~~~~~~~~~~
+
+This section describes the mapping of LLVM memory model onto AMDGPU machine code
+(see :ref:`memmodel`). *The implementation is WIP.*
+
+.. TODO
+ Update when implementation complete.
+
+ Support more relaxed OpenCL memory model to be controled by environment
+ component of target triple.
+
+The AMDGPU backend supports the memory synchronization scopes specified in
+:ref:`amdgpu-memory-scopes`.
+
+The code sequences used to implement the memory model are defined in table
+:ref:`amdgpu-amdhsa-memory-model-code-sequences-gfx6-gfx9-table`.
+
+The sequences specify the order of instructions that a single thread must
+execute. The ``s_waitcnt`` and ``buffer_wbinvl1_vol`` are defined with respect
+to other memory instructions executed by the same thread. This allows them to be
+moved earlier or later which can allow them to be combined with other instances
+of the same instruction, or hoisted/sunk out of loops to improve
+performance. Only the instructions related to the memory model are given;
+additional ``s_waitcnt`` instructions are required to ensure registers are
+defined before being used. These may be able to be combined with the memory
+model ``s_waitcnt`` instructions as described above.
+
+The AMDGPU memory model supports both the HSA [HSA]_ memory model, and the
+OpenCL [OpenCL]_ memory model. The HSA memory model uses a single happens-before
+relation for all address spaces (see :ref:`amdgpu-address-spaces`). The OpenCL
+memory model which has separate happens-before relations for the global and
+local address spaces, and only a fence specifying both global and local address
+space joins the relationships. Since the LLVM ``memfence`` instruction does not
+allow an address space to be specified the OpenCL fence has to convervatively
+assume both local and global address space was specified. However, optimizations
+can often be done to eliminate the additional ``s_waitcnt``instructions when
+there are no intervening corresponding ``ds/flat_load/store/atomic`` memory
+instructions. The code sequences in the table indicate what can be omitted for
+the OpenCL memory. The target triple environment is used to determine if the
+source language is OpenCL (see :ref:`amdgpu-opencl`).
+
+``ds/flat_load/store/atomic`` instructions to local memory are termed LDS
+operations.
+
+``buffer/global/flat_load/store/atomic`` instructions to global memory are
+termed vector memory operations.
+
+For GFX6-GFX9:
+
+* Each agent has multiple compute units (CU).
+* Each CU has multiple SIMDs that execute wavefronts.
+* The wavefronts for a single work-group are executed in the same CU but may be
+ executed by different SIMDs.
+* Each CU has a single LDS memory shared by the wavefronts of the work-groups
+ executing on it.
+* All LDS operations of a CU are performed as wavefront wide operations in a
+ global order and involve no caching. Completion is reported to a wavefront in
+ execution order.
+* The LDS memory has multiple request queues shared by the SIMDs of a
+ CU. Therefore, the LDS operations performed by different waves of a work-group
+ can be reordered relative to each other, which can result in reordering the
+ visibility of vector memory operations with respect to LDS operations of other
+ wavefronts in the same work-group. A ``s_waitcnt lgkmcnt(0)`` is required to
+ ensure synchonization between LDS operations and vector memory operations
+ between waves of a work-group, but not between operations performed by the
+ same wavefront.
+* The vector memory operations are performed as wavefront wide operations and
+ completion is reported to a wavefront in execution order. The exception is
+ that for GFX7-9 ``flat_load/store/atomic`` instructions can report out of
+ vector memory order if they access LDS memory, and out of LDS operation order
+ if they access global memory.
+* The vector memory operations access a vector L1 cache shared by all wavefronts
+ on a CU. Therefore, no special action is required for coherence between
+ wavefronts in the same work-group. A ``buffer_wbinvl1_vol`` is required for
+ coherence between waves executing in different work-groups as they may be
+ executing on different CUs.
+* The scalar memory operations access a scalar L1 cache shared by all wavefronts
+ on a group of CUs. The scalar and vector L1 caches are not coherent. However,
+ scalar operations are used in a restricted way so do not impact the memory
+ model. See :ref:`amdgpu-amdhsa-memory-spaces`.
+* The vector and scalar memory operations use an L2 cache shared by all CUs on
+ the same agent.
+* The L2 cache has independent channels to service disjoint ranges of virtual
+ addresses.
+* Each CU has a separate request queue per channel. Therefore, the vector and
+ scalar memory operations performed by waves executing in different work-groups
+ (which may be executing on different CUs) of an agent can be reordered
+ relative to each other. A ``s_waitcnt vmcnt(0)`` is required to ensure
+ synchonization between vector memory operations of different CUs. It ensures a
+ previous vector memory operation has completed before executing a subsequent
+ vector memory or LDS operation and so can be used to meet the requirements of
+ acquire and release.
+* The L2 cache can be kept coherent with other agents on some targets, or ranges
+ of virtual addresses can be set up to bypass it to ensure system coherence.
+
+Private address space uses ``buffer_load/store`` using the scratch V# (GFX6-8),
+or ``scratch_load/store`` (GFX9). Since only a single thread is accessing the
+memory, atomic memory orderings are not meaningful and all accesses are treated
+as non-atomic.
+
+Constant address space uses ``buffer/global_load`` instructions (or equivalent
+scalar memory instructions). Since the constant address space contents do not
+change during the execution of a kernel dispatch it is not legal to perform
+stores, and atomic memory orderings are not meaningful and all access are
+treated as non-atomic.
+
+A memory synchronization scope wider than work-group is not meaningful for the
+group (LDS) address space and is treated as work-group.
+
+The memory model does not support the region address space which is treated as
+non-atomic.
+
+Acquire memory ordering is not meaningful on store atomic instructions and is
+treated as non-atomic.
+
+Release memory ordering is not meaningful on load atomic instructions and is
+treated a non-atomic.
+
+Acquire-release memory ordering is not meaningful on load or store atomic
+instructions and is treated as acquire and release respectively.
+
+AMDGPU backend only uses scalar memory operations to access memory that is
+proven to not change during the execution of the kernel dispatch. This includes
+constant address space and global address space for program scope const
+variables. Therefore the kernel machine code does not have to maintain the
+scalar L1 cache to ensure it is coherent with the vector L1 cache. The scalar
+and vector L1 caches are invalidated between kernel dispatches by CP since
+constant address space data may change between kernel dispatch executions. See
+:ref:`amdgpu-amdhsa-memory-spaces`.
+
+The one exeception is if scalar writes are used to spill SGPR registers. In this
+case the AMDGPU backend ensures the memory location used to spill is never
+accessed by vector memory operations at the same time. If scalar writes are used
+then a ``s_dcache_wb`` is inserted before the ``s_endpgm`` and before a function
+return since the locations may be used for vector memory instructions by a
+future wave that uses the same scratch area, or a function call that creates a
+frame at the same address, respectively. There is no need for a ``s_dcache_inv``
+as all scalar writes are write-before-read in the same thread.
+
+Scratch backing memory (which is used for the private address space) is accessed
+with MTYPE NC_NV (non-coherenent non-volatile). Since the private address space
+is only accessed by a single thread, and is always write-before-read,
+there is never a need to invalidate these entries from the L1 cache. Hence all
+cache invalidates are done as ``*_vol`` to only invalidate the volatile cache
+lines.
+
+On dGPU the kernarg backing memory is accessed as UC (uncached) to avoid needing
+to invalidate the L2 cache. This also causes it to be treated as non-volatile
+and so is not invalidated by ``*_vol``. On APU it is accessed as CC (cache
+coherent) and so the L2 cache will coherent with the CPU and other agents.
+
+ .. table:: AMDHSA Memory Model Code Sequences GFX6-GFX9
+ :name: amdgpu-amdhsa-memory-model-code-sequences-gfx6-gfx9-table
+
+ ============ ============ ============== ========== =======================
+ LLVM Instr LLVM Memory LLVM Memory AMDGPU AMDGPU Machine Code
+ Ordering Sync Scope Address
+ Space
+ ============ ============ ============== ========== =======================
+ **Non-Atomic**
+ ---------------------------------------------------------------------------
+ load *none* *none* - global non-volatile
+ - generic 1. buffer/global/flat_load
+ volatile
+ 1. buffer/global/flat_load
+ glc=1
+ load *none* *none* - local 1. ds_load
+ store *none* *none* - global 1. buffer/global/flat_store
+ - generic
+ store *none* *none* - local 1. ds_store
+ **Unordered Atomic**
+ ---------------------------------------------------------------------------
+ load atomic unordered *any* *any* *Same as non-atomic*.
+ store atomic unordered *any* *any* *Same as non-atomic*.
+ atomicrmw unordered *any* *any* *Same as monotonic
+ atomic*.
+ **Monotonic Atomic**
+ ---------------------------------------------------------------------------
+ load atomic monotonic - singlethread - global 1. buffer/global/flat_load
+ - wavefront - generic
+ - workgroup
+ load atomic monotonic - singlethread - local 1. ds_load
+ - wavefront
+ - workgroup
+ load atomic monotonic - agent - global 1. buffer/global/flat_load
+ - system - generic glc=1
+ store atomic monotonic - singlethread - global 1. buffer/global/flat_store
+ - wavefront - generic
+ - workgroup
+ - agent
+ - system
+ store atomic monotonic - singlethread - local 1. ds_store
+ - wavefront
+ - workgroup
+ atomicrmw monotonic - singlethread - global 1. buffer/global/flat_atomic
+ - wavefront - generic
+ - workgroup
+ - agent
+ - system
+ atomicrmw monotonic - singlethread - local 1. ds_atomic
+ - wavefront
+ - workgroup
+ **Acquire Atomic**
+ ---------------------------------------------------------------------------
+ load atomic acquire - singlethread - global 1. buffer/global/ds/flat_load
+ - wavefront - local
+ - generic
+ load atomic acquire - workgroup - global 1. buffer/global_load
+ load atomic acquire - workgroup - local 1. ds/flat_load
+ - generic 2. s_waitcnt lgkmcnt(0)
+
+ - If OpenCL, omit
+ waitcnt.
+ - Must happen before
+ any following
+ global/generic
+ load/load
+ atomic/store/store
+ atomic/atomicrmw.
+ - Ensures any
+ following global
+ data read is no
+ older than the load
+ atomic value being
+ acquired.
+
+ load atomic acquire - agent - global 1. buffer/global_load
+ - system glc=1
+ 2. s_waitcnt vmcnt(0)
+
+ - Must happen before
+ following
+ buffer_wbinvl1_vol.
+ - Ensures the load
+ has completed
+ before invalidating
+ the cache.
+
+ 3. buffer_wbinvl1_vol
+
+ - Must happen before
+ any following
+ global/generic
+ load/load
+ atomic/atomicrmw.
+ - Ensures that
+ following
+ loads will not see
+ stale global data.
+
+ load atomic acquire - agent - generic 1. flat_load glc=1
+ - system 2. s_waitcnt vmcnt(0) &
+ lgkmcnt(0)
+
+ - If OpenCL omit
+ lgkmcnt(0).
+ - Must happen before
+ following
+ buffer_wbinvl1_vol.
+ - Ensures the flat_load
+ has completed
+ before invalidating
+ the cache.
+
+ 3. buffer_wbinvl1_vol
+
+ - Must happen before
+ any following
+ global/generic
+ load/load
+ atomic/atomicrmw.
+ - Ensures that
+ following loads
+ will not see stale
+ global data.
+
+ atomicrmw acquire - singlethread - global 1. buffer/global/ds/flat_atomic
+ - wavefront - local
+ - generic
+ atomicrmw acquire - workgroup - global 1. buffer/global_atomic
+ atomicrmw acquire - workgroup - local 1. ds/flat_atomic
+ - generic 2. waitcnt lgkmcnt(0)
+
+ - If OpenCL, omit
+ waitcnt.
+ - Must happen before
+ any following
+ global/generic
+ load/load
+ atomic/store/store
+ atomic/atomicrmw.
+ - Ensures any
+ following global
+ data read is no
+ older than the
+ atomicrmw value
+ being acquired.
+
+ atomicrmw acquire - agent - global 1. buffer/global_atomic
+ - system 2. s_waitcnt vmcnt(0)
+
+ - Must happen before
+ following
+ buffer_wbinvl1_vol.
+ - Ensures the
+ atomicrmw has
+ completed before
+ invalidating the
+ cache.
+
+ 3. buffer_wbinvl1_vol
+
+ - Must happen before
+ any following
+ global/generic
+ load/load
+ atomic/atomicrmw.
+ - Ensures that
+ following loads
+ will not see stale
+ global data.
+
+ atomicrmw acquire - agent - generic 1. flat_atomic
+ - system 2. s_waitcnt vmcnt(0) &
+ lgkmcnt(0)
+
+ - If OpenCL, omit
+ lgkmcnt(0).
+ - Must happen before
+ following
+ buffer_wbinvl1_vol.
+ - Ensures the
+ atomicrmw has
+ completed before
+ invalidating the
+ cache.
+
+ 3. buffer_wbinvl1_vol
+
+ - Must happen before
+ any following
+ global/generic
+ load/load
+ atomic/atomicrmw.
+ - Ensures that
+ following loads
+ will not see stale
+ global data.
+
+ fence acquire - singlethread *none* *none*
+ - wavefront
+ fence acquire - workgroup *none* 1. s_waitcnt lgkmcnt(0)
+
+ - If OpenCL and
+ address space is
+ not generic, omit
+ waitcnt. However,
+ since LLVM
+ currently has no
+ address space on
+ the fence need to
+ conservatively
+ always generate. If
+ fence had an
+ address space then
+ set to address
+ space of OpenCL
+ fence flag, or to
+ generic if both
+ local and global
+ flags are
+ specified.
+ - Must happen after
+ any preceding
+ local/generic load
+ atomic/atomicrmw
+ with an equal or
+ wider sync scope
+ and memory ordering
+ stronger than
+ unordered (this is
+ termed the
+ fence-paired-atomic).
+ - Must happen before
+ any following
+ global/generic
+ load/load
+ atomic/store/store
+ atomic/atomicrmw.
+ - Ensures any
+ following global
+ data read is no
+ older than the
+ value read by the
+ fence-paired-atomic.
+
+ fence acquire - agent *none* 1. s_waitcnt vmcnt(0) &
+ - system lgkmcnt(0)
+
+ - If OpenCL and
+ address space is
+ not generic, omit
+ lgkmcnt(0).
+ However, since LLVM
+ currently has no
+ address space on
+ the fence need to
+ conservatively
+ always generate
+ (see comment for
+ previous fence).
+ - Could be split into
+ separate s_waitcnt
+ vmcnt(0) and
+ s_waitcnt
+ lgkmcnt(0) to allow
+ them to be
+ independently moved
+ according to the
+ following rules.
+ - s_waitcnt vmcnt(0)
+ must happen after
+ any preceding
+ global/generic load
+ atomic/atomicrmw
+ with an equal or
+ wider sync scope
+ and memory ordering
+ stronger than
+ unordered (this is
+ termed the
+ fence-paired-atomic).
+ - s_waitcnt lgkmcnt(0)
+ must happen after
+ any preceding
+ group/generic load
+ atomic/atomicrmw
+ with an equal or
+ wider sync scope
+ and memory ordering
+ stronger than
+ unordered (this is
+ termed the
+ fence-paired-atomic).
+ - Must happen before
+ the following
+ buffer_wbinvl1_vol.
+ - Ensures that the
+ fence-paired atomic
+ has completed
+ before invalidating
+ the
+ cache. Therefore
+ any following
+ locations read must
+ be no older than
+ the value read by
+ the
+ fence-paired-atomic.
+
+ 2. buffer_wbinvl1_vol
+
+ - Must happen before
+ any following global/generic
+ load/load
+ atomic/store/store
+ atomic/atomicrmw.
+ - Ensures that
+ following loads
+ will not see stale
+ global data.
+
+ **Release Atomic**
+ ---------------------------------------------------------------------------
+ store atomic release - singlethread - global 1. buffer/global/ds/flat_store
+ - wavefront - local
+ - generic
+ store atomic release - workgroup - global 1. s_waitcnt lgkmcnt(0)
+ - generic
+ - If OpenCL, omit
+