summaryrefslogtreecommitdiffstats
path: root/llvm/lib/Target/NVPTX
Commit message (Collapse)AuthorAgeFilesLines
...
* [NVPTX] Renamed NVPTXLowerKernelArgs -> NVPTXLowerArgs. NFC.Artem Belevich2016-07-204-21/+21
| | | | | | | | After r276153 the pass applies to both kernels and regular functions. Differential Revision: https://reviews.llvm.org/D22583 llvm-svn: 276189
* [NVPTX] deal with all aggregate return types.Artem Belevich2016-07-201-6/+6
| | | | | | | | Fixes a crash in llvm_unreachable when a function has array return type. Differential Revision: https://reviews.llvm.org/D22524 llvm-svn: 276154
* [NVPTX] Improve lowering of byval args of device functions.Artem Belevich2016-07-203-32/+45
| | | | | | | | | | | | Avoid unnecessary spills of byval arguments of device functions to local space on SASS level and subsequent pointer conversion to generic address space that follows. Instead, make a local copy in IR, provide a way to access arguments directly, and let LLVM optimize the copy away when possible. Differential Review: https://reviews.llvm.org/D21421 llvm-svn: 276153
* [NVPTX] Make sure we adjust alignment at all call sitesArtem Belevich2016-07-181-2/+1
| | | | | | | .. including calls from kernel functions that were ignored by mistake before. llvm-svn: 275920
* [NVPTX] Force minimum alignment of 4 for byval arguments of device-side ↵Artem Belevich2016-07-182-6/+23
| | | | | | | | | | | | | | | | functions. Taking address of a byval variable in PTX is legal, but currently runs into miscompilation by ptxas on sm_50+ (NVIDIA issue 1789042). Work around the issue by enforcing minimum alignment on byval arguments of device functions. The change is a no-op on SASS level for sm_3x where ptxas already aligns local copy by at least 4. Differential Revision: https://reviews.llvm.org/D22428 llvm-svn: 275893
* [SelectionDAG] Get rid of bool parameters in SelectionDAG::getLoad, ↵Justin Lebar2016-07-151-32/+23
| | | | | | | | | | | | | | | | | | | | | | | getStore, and friends. Summary: Instead, we take a single flags arg (a bitset). Also add a default 0 alignment, and change the order of arguments so the alignment comes before the flags. This greatly simplifies many callsites, and fixes a bug in AMDGPUISelLowering, wherein the order of the args to getLoad was inverted. It also greatly simplifies the process of adding another flag to getLoad. Reviewers: chandlerc, tstellarAMD Subscribers: jholewinski, arsenm, jyknight, dsanders, nemanjai, llvm-commits Differential Revision: http://reviews.llvm.org/D22249 llvm-svn: 275592
* Rename AnalyzeBranch* to analyzeBranch*.Jacques Pienaar2016-07-152-6/+9
| | | | | | | | | | | | Summary: NFC. Rename AnalyzeBranch/AnalyzeBranchPredicate to analyzeBranch/analyzeBranchPredicate to follow LLVM coding style and be consistent with TargetInstrInfo's analyzeCompare and analyzeSelect. Reviewers: tstellarAMD, mcrosier Subscribers: mcrosier, jholewinski, jfb, arsenm, dschuff, jyknight, dsanders, nemanjai Differential Revision: https://reviews.llvm.org/D22409 llvm-svn: 275564
* NVPTX: Avoid implicit iterator conversions, NFCDuncan P. N. Exon Smith2016-07-083-22/+21
| | | | | | | | | | | | | | | | | Avoid implicit conversions from MachineInstrBundleIterator to MachineInstr* in the NVPTX backend, mainly by preferring MachineInstr& over MachineInstr* when a pointer isn't nullable and using range-based for loops. There was one piece of questionable code in NVPTXInstrInfo::AnalyzeBranch, where a condition checked a pointer converted from an iterator for nullptr. Since this case is impossible (moreover, the code above guarantees that the iterator is valid), I removed the check when I changed the pointer to a reference. Despite that case, there should be no functionality change here. llvm-svn: 274931
* NVPTX: Remove the legacy ptx intrinsicsJustin Bogner2016-07-073-136/+86
| | | | | | | | | | | | - Rename the ptx.read.* intrinsics to nvvm.read.ptx.sreg.* - some but not all of these registers were already accessible via the nvvm name. - Rename ptx.bar.sync nvvm.bar.sync, to match nvvm.bar0. There's a fair amount of code motion here, but it's all very mechanical. llvm-svn: 274769
* [NVPTX] Add sm_60, sm_61, sm_62 targets to LLVM.Justin Lebar2016-07-061-1/+13
| | | | | | | | | | Reviewers: tra Subscribers: jholewinski, llvm-commits Differential Revision: http://reviews.llvm.org/D22068 llvm-svn: 274674
* NVPTX: Replace uses of cuda.syncthreads with nvvm.barrier0Justin Bogner2016-07-062-4/+1
| | | | | | | Everywhere where cuda.syncthreads or __syncthreads is used, use the properly namespaced nvvm.barrier0 instead. llvm-svn: 274664
* NVPTX: Make the llvm.nvvm.shfl intrinsics and builtin names consistentJustin Bogner2016-07-061-8/+8
| | | | | | | The intrinsics here use nvvm, but the builtins and tablegen variable names were using ptx. Stick to the modern names here. llvm-svn: 274662
* Use arrays or initializer lists to feed ArrayRefs instead of SmallVector ↵Benjamin Kramer2016-07-021-4/+1
| | | | | | | | where possible. No functionality change intended. llvm-svn: 274431
* Delete MCCodeGenInfo.Rafael Espindola2016-06-301-14/+0
| | | | | | | MC doesn't really care about CodeGen stuff, so this was just complicating target initialization. llvm-svn: 274258
* Revert r273313 "[NVPTX] Improve lowering of byval args of device functions."Artem Belevich2016-06-293-79/+15
| | | | | | The change causes llvm crash in some unoptimized builds. llvm-svn: 274163
* Only emit extension for zeroext/signext arguments if type is < 32 bitsJustin Holewinski2016-06-271-2/+2
| | | | | | | | | | Reviewers: jingyue, jlebar Subscribers: jholewinski Differential Revision: http://reviews.llvm.org/D21756 llvm-svn: 273922
* [NVPTX] Improve lowering of byval args of device functions.Artem Belevich2016-06-213-15/+79
| | | | | | | | | | | | | | | | | | | | | Avoid unnecessary spills of such vars to local space on SASS level and pointer space conversion. Instead, make a local copy with appropriate addrspacecasts and let LLVM optimize them away when possible. This allows loading value of the argument using [symbol+offset] instead of converting argument to general space pointer and using it for indexing (which also implicitly converts param space pointer to local space one on SASS level and triggers copying of argument into local space in the process). This reduces call overhead, uses less registers and reduces overall SASS size by 2-4%. Differential Review: http://reviews.llvm.org/D21421 llvm-svn: 273313
* Pass DebugLoc and SDLoc by const ref.Benjamin Kramer2016-06-125-26/+28
| | | | | | | | This used to be free, copying and moving DebugLocs became expensive after the metadata rewrite. Passing by reference eliminates a ton of track/untrack operations. No functionality change intended. llvm-svn: 272512
* [NVPTX] Add intrinsics for shfl instructions.Justin Lebar2016-06-091-1/+42
| | | | | | | | | | | | | | | Summary: Currently clang emits these instructions via inline (volatile) asm in the CUDA headers. Switching to intrinsics will let the optimizer reason across calls to these intrinsics. Reviewers: tra Subscribers: llvm-commits, jholewinski Differential Revision: http://reviews.llvm.org/D21160 llvm-svn: 272298
* Apply most suggestions of clang-tidy's performance-unnecessary-value-paramBenjamin Kramer2016-06-086-11/+12
| | | | | | | Avoids unnecessary copies. All changes audited & pass tests with asan. No functional change intended. llvm-svn: 272190
* Avoid copies of std::strings and APInt/APFloats where we only read from itBenjamin Kramer2016-06-082-2/+2
| | | | | | | | As suggested by clang-tidy's performance-unnecessary-copy-initialization. This can easily hit lifetime issues, so I audited every change and ran the tests under asan, which came back clean. llvm-svn: 272126
* Apply clang-tidy's misc-move-constructor-init throughout LLVM.Benjamin Kramer2016-05-271-1/+2
| | | | | | No functionality change intended, maybe a tiny performance improvement. llvm-svn: 270997
* Avoid some copies by using const references.Benjamin Kramer2016-05-271-1/+1
| | | | | | | clang-tidy's performance-unnecessary-copy-initialization with some manual fixes. No functional changes intended. llvm-svn: 270988
* Init member structs in constructor.Artem Belevich2016-05-261-3/+9
| | | | | | | Fixes build error on windows where MSVC does not support list initialization inside member initializer list. llvm-svn: 270877
* [NVPTX] Added NVVMIntrRange passArtem Belevich2016-05-264-0/+159
| | | | | | | | | | | | NVVMIntrRange adds !range metadata to calls of NVVM intrinsics that return values within known limited range. This allows LLVM to generate optimal code for indexing arrays based on tid/ctaid which is a frequently used pattern in CUDA code. Differential Revision: http://reviews.llvm.org/D20644 llvm-svn: 270872
* [NVPTX] Don't (incorrectly) say that the NVVMReflect pass preserves all ↵Justin Lebar2016-05-251-3/+0
| | | | | | | | | | | | analyses. Reviewers: tra Subscribers: jholewinski, llvm-commits Differential Revision: http://reviews.llvm.org/D20585 llvm-svn: 270790
* Delete Reloc::Default.Rafael Espindola2016-05-183-13/+17
| | | | | | | | | | | | Having an enum member named Default is quite confusing: Is it distinct from the others? This patch removes that member and instead uses Optional<Reloc> in places where we have a user input that still hasn't been maped to the default value, which is now clear has no be one of the remaining 3 options. llvm-svn: 269988
* SDAG: Implement Select instead of SelectImpl in NVPTXDAGToDAGISelJustin Bogner2016-05-132-213/+236
| | | | | | | | | | - Where we were returning a node before, call ReplaceNode instead. - Where we would return null to fall back to another selector, rename the method to try* and return a bool for success. Part of llvm.org/pr26808. llvm-svn: 269483
* Return a StringRef from getSection.Rafael Espindola2016-05-111-1/+1
| | | | | | This is similar to how getName is handled. llvm-svn: 269218
* CodeGen: Move TargetPassConfig from Passes.h to an own header; NFCMatthias Braun2016-05-101-0/+1
| | | | | | | | Many files include Passes.h but only a fraction needs to know about the TargetPassConfig class. Move it into an own header. Also rename Passes.cpp to TargetPassConfig.cpp while we are at it. llvm-svn: 269011
* [NVPTX] Change begin/end inline asm comments to "begin/end inline asm".Justin Lebar2016-05-101-2/+2
| | | | | | | Previously it was just "// inline asm", which made it tricky to read code with lots of inline assembly. llvm-svn: 268994
* SDAG: Rename Select->SelectImpl and repurpose Select as returning voidJustin Bogner2016-05-052-3/+3
| | | | | | | | | | | | | | This is a step towards removing the rampant undefined behaviour in SelectionDAG, which is a part of llvm.org/PR26808. We rename SelectionDAGISel::Select to SelectImpl and update targets to match, and then change Select to return void and consolidate the sketchy behaviour we're trying to get away from there. Next, we'll update backends to implement `void Select(...)` instead of SelectImpl and eventually drop the base Select implementation. llvm-svn: 268693
* [NVPTX] Fix sign/zero-extending ldg/ldu instruction selectionJustin Holewinski2016-05-023-48/+77
| | | | | | | | | | | | | | | | | Summary: We don't have sign-/zero-extending ldg/ldu instructions defined, so we need to emulate them with explicit CVTs. We were originally handling the i8 case, but not any other cases. Fixes PR26185 Reviewers: jingyue, jlebar Subscribers: jholewinski Differential Revision: http://reviews.llvm.org/D19615 llvm-svn: 268272
* [CodeGen] Default CTTZ_ZERO_UNDEF/CTLZ_ZERO_UNDEF to Expand in ↵Craig Topper2016-04-281-6/+0
| | | | | | TargetLoweringBase. This is what the majority of the targets want and removes a bunch of code. Set it to Legal explicitly in the few cases where that's the desired behavior. llvm-svn: 267853
* [NVPTX] Run NVVMReflect at the beginning of IR passes.Justin Lebar2016-04-272-0/+10
| | | | | | | | | | | | | | | | | | | | | | | | | Summary: Currently the NVVMReflect pass is run at the beginning of our backend passes. But really, it should be run as early as possible, as it's simply resolving an "if" statement in code. So copy it into TargetMachine::addEarlyAsPossiblePasses. We still run it at the beginning of the backend passes, since it's needed for correctness when lowering to nvptx. (Specifically, NVVMReflect changes each call to the __nvvm_reflect function or llvm.nvvm.reflect intrinsic into an integer constant, based on the pass's configuration. Clearly we miss many optimization opportunities if we perform this transformation at the beginning of codegen.) Reviewers: rnk Subscribers: tra, llvm-commits, jholewinski Differential Revision: http://reviews.llvm.org/D18616 llvm-svn: 267765
* Add optimization bisect opt-in calls for NVPTX passesAndrew Kaylor2016-04-265-1/+13
| | | | | | Differential Revision: http://reviews.llvm.org/D19518 llvm-svn: 267635
* [NVPTX] Fix some usages of CodeGenOpt::None.Jingyue Wu2016-04-261-5/+9
| | | | | | | | | | NVPTXLowerKernelArgs is required for correctness, so it should not be guarded by CodeGenOpt::None. NVPTXPeephole is optimization only, so it should be skipped when CodeGenOpt::None. llvm-svn: 267619
* [CodeGen] Add getBuildVector and getSplatBuildVector helpers. NFCI.Ahmed Bougacha2016-04-261-3/+3
| | | | | | Differential Revision: http://reviews.llvm.org/D17176 llvm-svn: 267606
* [NVPTX] Set ctlz_zero_undef to Expand so LegalizeDAG will convert it to ↵Craig Topper2016-04-232-10/+3
| | | | | | ctlz. Remove the now unneccessary isel patterns. NFC llvm-svn: 267265
* Disable the PatchableFunction pass for NVPTX & WasmSanjoy Das2016-04-191-0/+1
| | | | | | | PatchableFunction requires AllVRegsAllocated that these targets don't provide. llvm-svn: 266720
* [NFC] Header cleanupMehdi Amini2016-04-186-10/+0
| | | | | | | | | | | | | | Removed some unused headers, replaced some headers with forward class declarations. Found using simple scripts like this one: clear && ack --cpp -l '#include "llvm/ADT/IndexedMap.h"' | xargs grep -L 'IndexedMap[<]' | xargs grep -n --color=auto 'IndexedMap' Patch by Eugene Kosov <claprix@yandex.ru> Differential Revision: http://reviews.llvm.org/D19219 From: Mehdi Amini <mehdi.amini@apple.com> llvm-svn: 266595
* [NVPTX] Set NVPTXTTI::getInliningThresholdMultiplier to 5.Justin Lebar2016-04-151-0/+4
| | | | | | | | | | | | | | | Summary: Calls on NVPTX are unusually expensive (for one thing, lots of state needs to be saved to memory, which is slow), so make the inlininer much more aggressive. Reviewers: chandlerc Subscribers: jholewinski, llvm-commits, tra Differential Revision: http://reviews.llvm.org/D18561 llvm-svn: 266406
* IR: Introduce ConstantAggregate, NFCDuncan P. N. Exon Smith2016-04-052-4/+2
| | | | | | | | | | | | Add a common parent class for ConstantArray, ConstantVector, and ConstantStruct called ConstantAggregate. These are the aggregate subclasses of Constant that take operands. This is mainly a cleanup, adding common `isa` target and removing duplicated code. However, it also simplifies caching which constants point transitively at `GlobalValue` (a possible future direction). llvm-svn: 265466
* [NVPTX] Handle ldg created from sign-/zero-extended loadJustin Holewinski2016-04-052-4/+91
| | | | | | | | | | Reviewers: jingyue Subscribers: jholewinski Differential Revision: http://reviews.llvm.org/D18053 llvm-svn: 265389
* [NVPTX] Add a truncate DAG node to some calls.Justin Lebar2016-04-011-2/+10
| | | | | | | | | | | | | | | | | | | | | | | Summary: Previously, we were running afoul of the assertion EVT(CLI.Ins[i].VT) == InVals[i].getValueType() && "LowerCall emitted a value with the wrong type!" in SelectionDAGBuilder.cpp when running the NVPTX/i8-param.ll test. This is because our backend (for some reason) treats small return values as i32, but it wasn't ever truncating the i32 back down to the expected width in the DAG. Unclear to me whether this fixes any actual bugs -- in this test, at least, the generated code is unchanged. Reviewers: jingyue Subscribers: llvm-commits, tra, jholewinski Differential Revision: http://reviews.llvm.org/D17872 llvm-svn: 265091
* [NVPTX] Read __CUDA_FTZ from module flags in NVVMReflect.Justin Lebar2016-04-011-7/+17
| | | | | | | | | | | | | | | | | | | Summary: Previously the NVVMReflect pass would read its configuration from command-line flags or a static configuration given to the pass at instantiation time. This doesn't quite work for clang's use-case. It needs to pass a value for __CUDA_FTZ down on a per-module basis. We use a module flag for this, so the NVVMReflect pass needs to be updated to read said flag. Reviewers: tra, rnk Subscribers: cfe-commits, jholewinski Differential Revision: http://reviews.llvm.org/D18672 llvm-svn: 265090
* [NVPTX] Annotate some instructions as hasSideEffects = 0.Justin Lebar2016-04-012-146/+171
| | | | | | | | | | | | | | | | | | | | Summary: Tablegen tries to infer this from the selection DAG patterns defined for the instructions, but it can't always. An instructive example is CLZr64. CLZr32 is correctly inferred to have no side-effects, but the selection DAG pattern for CLZr64 is slightly more complicated, and in particular the ctlz DAG node is not at the root of the pattern. Thus tablegen can't infer that CLZr64 has no side-effects. Reviewers: jholewinski Subscribers: jholewinski, tra, llvm-commits Differential Revision: http://reviews.llvm.org/D17472 llvm-svn: 265089
* Change eliminateCallFramePseudoInstr() to return an iteratorHans Wennborg2016-03-312-3/+3
| | | | | | | | | | | | | | | | | | | | | This will become necessary in a subsequent change to make this method merge adjacent stack adjustments, i.e. it might erase the previous and/or next instruction. It also greatly simplifies the calls to this function from Prolog- EpilogInserter. Previously, that had a bunch of logic to resume iteration after the call; now it just continues with the returned iterator. Note that this changes the behaviour of PEI a little. Previously, it attempted to re-visit the new instruction created by eliminateCallFramePseudoInstr(). That code was added in r36625, but I can't see any reason for it: the new instructions will obviously not be pseudo instructions, they will not have FrameIndex operands, and we have already accounted for the stack adjustment. Differential Revision: http://reviews.llvm.org/D18627 llvm-svn: 265036
* [NVPTX] Make NVVMReflect a function pass.Justin Lebar2016-03-302-102/+69
| | | | | | | | | | | | | | | Summary: Currently it's a module pass. Make it a function pass so that we can move it to PassManagerBuilder's EP_EarlyAsPossible extension point, which only accepts function passes. Reviewers: rnk Subscribers: tra, llvm-commits, jholewinski Differential Revision: http://reviews.llvm.org/D18615 llvm-svn: 264919
* [NVPTX] Avoid temporary std::string and make single-use function local to ↵Benjamin Kramer2016-03-302-5/+4
| | | | | | | | the cpp file. No functionality change intended. llvm-svn: 264861
OpenPOWER on IntegriCloud