summaryrefslogtreecommitdiffstats
path: root/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
Commit message (Collapse)AuthorAgeFilesLines
* NVPTX: Remove implicit ilist iterator conversions, NFCDuncan P. N. Exon Smith2015-10-201-4/+3
| | | | llvm-svn: 250779
* [NVPTX] Remove dead code.Benjamin Kramer2015-10-151-14/+0
| | | | | | I left helpers that look useful for debugging alone. NFC. llvm-svn: 250410
* [NVPTX] Allow undef value as global initializerJingyue Wu2015-08-221-3/+5
| | | | | | | | | | | | | | | | | | Summary: __shared__ variable may now emit undef value as initializer, do not throw error on that. Test Plan: test/CodeGen/NVPTX/global-addrspace.ll Patch by Xuetian Weng Reviewers: jholewinski, tra, jingyue Subscribers: llvm-commits, jholewinski Differential Revision: http://reviews.llvm.org/D12242 llvm-svn: 245785
* De-constify pointers to Type since they can't be modified. NFCCraig Topper2015-08-011-14/+14
| | | | | | This was already done in most places a while ago. This just fixes the ones that crept in over time. llvm-svn: 243842
* Refactor: Simplify boolean conditional return statements in lib/Target/NVPTXJingyue Wu2015-07-311-9/+4
| | | | | | | | | | | | Summary: Use clang-tidy to simplify boolean conditional return statements Reviewers: rafael, echristo, chandlerc, bkramer, craig.topper, dexonsmith, chapuni, eliben, jingyue, jholewinski Subscribers: llvm-commits, jholewinski Differential Revision: http://reviews.llvm.org/D9983 llvm-svn: 243734
* Move most user of TargetMachine::getDataLayout to the Module oneMehdi Amini2015-07-161-36/+36
| | | | | | | | | | | | | | | | | | | | | Summary: This change is part of a series of commits dedicated to have a single DataLayout during compilation by using always the one owned by the module. This patch is quite boring overall, except for some uglyness in ASMPrinter which has a getDataLayout function but has some clients that use it without a Module (llmv-dsymutil, llvm-dwarfdump), so some methods are taking a DataLayout as parameter. Reviewers: echristo Subscribers: yaron.keren, rafael, llvm-commits, jholewinski Differential Revision: http://reviews.llvm.org/D11090 From: Mehdi Amini <mehdi.amini@apple.com> llvm-svn: 242386
* Make TargetLowering::getPointerTy() taking DataLayout as an argumentMehdi Amini2015-07-091-2/+2
| | | | | | | | | | | | | | | | Summary: This change is part of a series of commits dedicated to have a single DataLayout during compilation by using always the one owned by the module. Reviewers: echristo Subscribers: jholewinski, ted, yaron.keren, rafael, llvm-commits Differential Revision: http://reviews.llvm.org/D11028 From: Mehdi Amini <mehdi.amini@apple.com> llvm-svn: 241775
* Redirect DataLayout from TargetMachine to Module in ComputeValueVTs()Mehdi Amini2015-07-091-5/+5
| | | | | | | | | | | | | | | | | | | | Summary: Avoid using the TargetMachine owned DataLayout and use the Module owned one instead. This requires passing the DataLayout up the stack to ComputeValueVTs(). This change is part of a series of commits dedicated to have a single DataLayout during compilation by using always the one owned by the module. Reviewers: echristo Subscribers: jholewinski, yaron.keren, rafael, llvm-commits Differential Revision: http://reviews.llvm.org/D11019 From: Mehdi Amini <mehdi.amini@apple.com> llvm-svn: 241773
* Simplify the Mangler interface now that DataLayout is mandatory.Rafael Espindola2015-06-231-1/+1
| | | | | | | We only need to pass in a DataLayout when mangling a raw string, not when constructing the mangler. llvm-svn: 240405
* Revert r240137 (Fixed/added namespace ending comments using clang-tidy. NFC)Alexander Kornienko2015-06-231-1/+1
| | | | | | Apparently, the style needs to be agreed upon first. llvm-svn: 240390
* Fixed/added namespace ending comments using clang-tidy. NFCAlexander Kornienko2015-06-191-1/+1
| | | | | | | | | | | | | The patch is generated using this command: tools/clang/tools/extra/clang-tidy/tool/run-clang-tidy.py -fix \ -checks=-*,llvm-namespace-comment -header-filter='llvm/.*|clang/.*' \ llvm/lib/ Thanks to Eugene Kosov for the original patch! llvm-svn: 240137
* Clean up redundant copies of Triple objects. NFCDaniel Sanders2015-06-161-2/+2
| | | | | | | | | | | | | | Summary: Reviewers: rengolin Reviewed By: rengolin Subscribers: llvm-commits, rengolin, jholewinski Differential Revision: http://reviews.llvm.org/D10382 llvm-svn: 239823
* Remove MachineModuleInfo::UsedFunctions as it has no users.Pete Cooper2015-06-111-1/+0
| | | | | | | | | | It hasn't been used since r130964. This also removes MachineModuleInfo::isUsedFunction and MachineModuleInfo::AnalyzeModule, both of which were only there to support UsedFunctions. llvm-svn: 239501
* Replace string GNU Triples with llvm::Triple in MCSubtargetInfo and ↵Daniel Sanders2015-06-101-1/+1
| | | | | | | | | | | | | | | | | | create*MCSubtargetInfo(). NFC. Summary: This continues the patch series to eliminate StringRef forms of GNU triples from the internals of LLVM that began in r239036. Reviewers: rafael Reviewed By: rafael Subscribers: rafael, ted, jfb, llvm-commits, rengolin, jholewinski Differential Revision: http://reviews.llvm.org/D10311 llvm-svn: 239467
* The constant initialization for globals in NVPTX is generated as an Samuel Antao2015-06-091-9/+33
| | | | | | | | | | array of bytes. The generation of this byte arrays was expecting the host to be little endian, which prevents big endian hosts to be used in the generation of the PTX code. This patch fixes the problem by changing the way the bytes are extracted so that it works for either little and big endian. llvm-svn: 239412
* MC: Add target hook to control symbol quotingMatt Arsenault2015-06-091-24/+40
| | | | llvm-svn: 239370
* Push constness through LoopInfo::isLoopHeader and clean it up a bit.Benjamin Kramer2015-06-021-2/+1
| | | | | | NFC. llvm-svn: 238843
* MC: Clean up MCExpr naming. NFC.Jim Grosbach2015-05-301-13/+13
| | | | llvm-svn: 238634
* Don't call utostr in Twine/raw_ostream contexts.Benjamin Kramer2015-05-281-6/+4
| | | | | | Creating temporary std::strings there is unnecessary. llvm-svn: 238412
* MC: Clean up method names in MCContext.Jim Grosbach2015-05-181-2/+2
| | | | | | | The naming was a mish-mash of old and new style. Update to be consistent with the new. NFC. llvm-svn: 237594
* Remove MCAssembler.h include from MCStreamer.h and fix users of MCStreamer.hPete Cooper2015-05-151-0/+1
| | | | llvm-svn: 237483
* MC: Modernize MCOperand API naming. NFC.Jim Grosbach2015-05-131-6/+6
| | | | | | MCOperand::Create*() methods renamed to MCOperand::create*(). llvm-svn: 237275
* IR: Give 'DI' prefix to debug info metadataDuncan P. N. Exon Smith2015-04-291-3/+3
| | | | | | | | | | | | | | | | | | | | | | | | | | | | Finish off PR23080 by renaming the debug info IR constructs from `MD*` to `DI*`. The last of the `DIDescriptor` classes were deleted in r235356, and the last of the related typedefs removed in r235413, so this has all baked for about a week. Note: If you have out-of-tree code (like a frontend), I recommend that you get everything compiling and tests passing with the *previous* commit before updating to this one. It'll be easier to keep track of what code is using the `DIDescriptor` hierarchy and what you've already updated, and I think you're extremely unlikely to insert bugs. YMMV of course. Back to *this* commit: I did this using the rename-md-di-nodes.sh upgrade script I've attached to PR23080 (both code and testcases) and filtered through clang-format-diff.py. I edited the tests for test/Assembler/invalid-generic-debug-node-*.ll by hand since the columns were off-by-three. It should work on your out-of-tree testcases (and code, if you've followed the advice in the previous paragraph). Some of the tests are in badly named files now (e.g., test/Assembler/invalid-mdcompositetype-missing-tag.ll should be 'dicompositetype'); I'll come back and move the files in a follow-up commit. llvm-svn: 236120
* [NVPTX] Handle addrspacecast constant expressions in aggregate initializersJustin Holewinski2015-04-281-0/+206
| | | | | | | | | | | We need to track if an AddrSpaceCast expression was seen when generating an MCExpr for a ConstantExpr. This change introduces a custom lowerConstant method to the NVPTX asm printer that will create NVPTXGenericMCSymbolRefExpr nodes at the appropriate places to encode the information that a given symbol needs to be casted to a generic address. llvm-svn: 236000
* [AsmPrinter] Make AsmPrinter's OutStreamer member a unique_ptr.Lang Hames2015-04-241-23/+23
| | | | | | | AsmPrinter owns the OutStreamer, so an owning pointer makes sense here. Using a reference for this is crufty. llvm-svn: 235752
* [NVPTX] Emits "generic()" depending on the original address spaceJingyue Wu2015-04-241-5/+5
| | | | | | | | | | | | | | | | | | Summary: Fixes a bug in the NVPTX codegen. The code used to miss necessary "generic()" on aggregates of addrspacecasts. Test Plan: addrspacecast-gvar.ll Reviewers: eliben, jholewinski Reviewed By: jholewinski Subscribers: jholewinski, llvm-commits Differential Revision: http://reviews.llvm.org/D9130 llvm-svn: 235689
* DebugInfo: Gut DIScope, DIEnumerator and DISubrangeDuncan P. N. Exon Smith2015-04-161-3/+3
| | | | | | The only class the still has API left is `DIDescriptor` itself. llvm-svn: 235067
* DebugInfo: Gut DICompileUnit and DIFileDuncan P. N. Exon Smith2015-04-151-3/+3
| | | | | | | Continuing gutting `DIDescriptor` subclasses; this edition, `DICompileUnit` and `DIFile`. In the name of PR23080. llvm-svn: 235055
* DebugInfo: Gut DISubprogram and DILexicalBlock*Duncan P. N. Exon Smith2015-04-141-3/+3
| | | | | | | Gut the `DIDescriptor` wrappers around `MDLocalScope` subclasses. Note that `DILexicalBlock` wraps `MDLexicalBlockBase`, not `MDLexicalBlock`. llvm-svn: 234850
* Reduce dyn_cast<> to isa<> or cast<> where possible.Benjamin Kramer2015-04-101-3/+2
| | | | | | No functional change intended. llvm-svn: 234586
* CodeGen: Stop using DIDescriptor::is*() and auto-castingDuncan P. N. Exon Smith2015-04-061-4/+1
| | | | | | Same as r234255, but for lib/CodeGen and lib/Target. llvm-svn: 234258
* CodeGen: Use the new DebugLoc API, NFCDuncan P. N. Exon Smith2015-03-301-7/+3
| | | | | | Update lib/CodeGen (and lib/Target) to use the new `DebugLoc` API. llvm-svn: 233582
* Remove more superfluous .str() and replace std::string concatenation with Twine.Yaron Keren2015-03-301-21/+22
| | | | | | Following r233392, http://llvm.org/viewvc/llvm-project?rev=233392&view=rev. llvm-svn: 233555
* Disabling warnings for MSVC build to enable /W4 use.Andrew Kaylor2015-03-241-2/+1
| | | | | | Differential Revision: http://reviews.llvm.org/D8572 llvm-svn: 233133
* Simplify boolean expressions with true and false using clang-tidyEli Bendersky2015-03-231-10/+10
| | | | | | | | Patch by Richard (legalize@xmission.com) Differential Revision: http://reviews.llvm.org/D8521 llvm-svn: 232961
* NVPTX: Remove copy of LLVMInitializeNVPTXAsmPrinter.Benjamin Kramer2015-03-101-7/+0
| | | | | | | | If anyone is using this for some strange reason, LLVMInitializeNVPTXAsmPrinter does exactly the same thing and is what other LLVM tools are calling. llvm-svn: 231810
* DataLayout is mandatory, update the API to reflect it with references.Mehdi Amini2015-03-101-2/+2
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | Summary: Now that the DataLayout is a mandatory part of the module, let's start cleaning the codebase. This patch is a first attempt at doing that. This patch is not exactly NFC as for instance some places were passing a nullptr instead of the DataLayout, possibly just because there was a default value on the DataLayout argument to many functions in the API. Even though it is not purely NFC, there is no change in the validation. I turned as many pointer to DataLayout to references, this helped figuring out all the places where a nullptr could come up. I had initially a local version of this patch broken into over 30 independant, commits but some later commit were cleaning the API and touching part of the code modified in the previous commits, so it seemed cleaner without the intermediate state. Test Plan: Reviewers: echristo Subscribers: llvm-commits From: Mehdi Amini <mehdi.amini@apple.com> llvm-svn: 231740
* Avoid using a self-referential initializer and fix up uses.Eric Christopher2015-02-191-3/+3
| | | | llvm-svn: 229790
* Remove all use of is64bit off of NVPTXSubtarget and clean up codeEric Christopher2015-02-191-2/+2
| | | | | | | accordingly. This changes the constructors of a number of classes that don't need to know the subtarget's 64-bitness. llvm-svn: 229787
* Remove all use of getDrvInterface off of NVPTXSubtarget and cleanEric Christopher2015-02-191-36/+10
| | | | | | | up code accordingly. Delete code that was checking for all cases of an enum. llvm-svn: 229786
* Migrate the NVPTX backend asm printer to a per function subtarget.Eric Christopher2015-02-191-34/+40
| | | | | | | | | | | This involved moving two non-subtarget dependent features (64-bitness and the driver interface) to the NVPTX target machine and updating the uses (or migrating around the subtarget use for ease of review). Otherwise use the cached subtarget or create a default subtarget based on the TargetMachine cpu and feature string for the module level assembler emission. llvm-svn: 229785
* Remove usernames from TODOs, NFCJingyue Wu2015-02-031-3/+2
| | | | | | making the style consistent with the rest llvm-svn: 227991
* Resurrect the assertion removed by r227717Jingyue Wu2015-02-021-2/+1
| | | | | | | | | | | | | | Summary: MSVC can compile "LoopID->getOperand(0) == LoopID" when LoopID is MDNode*. Test Plan: no regression Reviewers: mkuper Subscribers: jholewinski, llvm-commits Differential Revision: http://reviews.llvm.org/D7327 llvm-svn: 227853
* [NVPTX] Emit .pragma "nounroll" for loops marked with nounrollJingyue Wu2015-02-011-0/+40
| | | | | | | | | | | | | | | | | | | | | | | Summary: CUDA driver can unroll loops when jit-compiling PTX. To prevent CUDA driver from unrolling a loop marked with llvm.loop.unroll.disable is not unrolled by CUDA driver, we need to emit .pragma "nounroll" at the header of that loop. This patch also extracts getting unroll metadata from loop ID metadata into a shared helper function. Test Plan: test/CodeGen/NVPTX/nounroll.ll Reviewers: eliben, meheff, jholewinski Reviewed By: jholewinski Subscribers: jholewinski, llvm-commits Differential Revision: http://reviews.llvm.org/D7041 llvm-svn: 227703
* Move DataLayout back to the TargetMachine from TargetSubtargetInfoEric Christopher2015-01-261-7/+7
| | | | | | | | | | | | | | | | | | | derived classes. Since global data alignment, layout, and mangling is often based on the DataLayout, move it to the TargetMachine. This ensures that global data is going to be layed out and mangled consistently if the subtarget changes on a per function basis. Prior to this all targets(*) have had subtarget dependent code moved out and onto the TargetMachine. *One target hasn't been migrated as part of this change: R600. The R600 port has, as a subtarget feature, the size of pointers and this affects global data layout. I've currently hacked in a FIXME to enable progress, but the port needs to be updated to either pass the 64-bitness to the TargetMachine, or fix the DataLayout to avoid subtarget dependent features. llvm-svn: 227113
* [cleanup] Re-sort all the #include lines in LLVM usingChandler Carruth2015-01-141-1/+1
| | | | | | | | | | | utils/sort_includes.py. I clearly haven't done this in a while, so more changed than usual. This even uncovered a missing include from the InstrProf library that I've added. No functionality changed here, just mechanical cleanup of the include order. llvm-svn: 225974
* Replace several 'assert(false' with 'llvm_unreachable' or fold a condition ↵Craig Topper2015-01-051-11/+9
| | | | | | into the assert. llvm-svn: 225160
* [NVPTX] Fix bugs related to isSingleValueTypeJingyue Wu2014-12-171-13/+2
| | | | | | | | | | | | | | | | | | | | Summary: With isSingleValueType starting to treat vector types as single-value types, code that uses this interface needs to be updated. Test Plan: vector-global.ll nvcl-param-align.ll Reviewers: jholewinski Reviewed By: jholewinski Subscribers: llvm-commits, meheff, eliben, jholewinski Differential Revision: http://reviews.llvm.org/D6573 llvm-svn: 224440
* NVPTX: Remove duplicate of AsmPrinter::lowerConstantMatt Arsenault2014-12-161-155/+1
| | | | llvm-svn: 224355
* NVPTX: Delete dead codeDuncan P. N. Exon Smith2014-12-031-5/+0
| | | | | | `MDNode` does not inherit from `User`, and it never has a name. llvm-svn: 223198
OpenPOWER on IntegriCloud