summaryrefslogtreecommitdiff
path: root/lib/Target/NVPTX
AgeCommit message (Collapse)Author
2017-12-28Avoid int to string conversion in Twine or raw_ostream contexts.Benjamin Kramer
Some output changes from uppercase hex to lowercase hex, no other functionality change intended. git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@321526 91177308-0d34-0410-b5e6-96231b3b80d8
2017-12-22(Re-landing) Expose a TargetMachine::getTargetTransformInfo functionSanjoy Das
Re-land r321234. It had to be reverted because it broke the shared library build. The shared library build broke because there was a missing LLVMBuild dependency from lib/Passes (which calls TargetMachine::getTargetIRAnalysis) to lib/Target. As far as I can tell, this problem was always there but was somehow masked before (perhaps because TargetMachine::getTargetIRAnalysis was a virtual function). Original commit message: This makes the TargetMachine interface a bit simpler. We still need the std::function in TargetIRAnalysis to avoid having to add a dependency from Analysis to Target. See discussion: http://lists.llvm.org/pipermail/llvm-dev/2017-December/119749.html I avoided adding all of the backend owners to this review since the change is simple, but let me know if you feel differently about this. Reviewers: echristo, MatzeB, hfinkel Reviewed By: hfinkel Subscribers: jholewinski, jfb, arsenm, dschuff, mcrosier, sdardis, nemanjai, nhaehnle, javed.absar, sbc100, jgravelle-google, aheejin, kbarton, llvm-commits Differential Revision: https://reviews.llvm.org/D41464 git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@321375 91177308-0d34-0410-b5e6-96231b3b80d8
2017-12-21Revert "Expose a TargetMachine::getTargetTransformInfo function"Sanjoy Das
This reverts commit r321234. It breaks the -DBUILD_SHARED_LIBS=ON build. git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@321243 91177308-0d34-0410-b5e6-96231b3b80d8
2017-12-21Expose a TargetMachine::getTargetTransformInfo functionSanjoy Das
Summary: This makes the TargetMachine interface a bit simpler. We still need the std::function in TargetIRAnalysis to avoid having to add a dependency from Analysis to Target. See discussion: http://lists.llvm.org/pipermail/llvm-dev/2017-December/119749.html I avoided adding all of the backend owners to this review since the change is simple, but let me know if you feel differently about this. Reviewers: echristo, MatzeB, hfinkel Reviewed By: hfinkel Subscribers: jholewinski, jfb, arsenm, dschuff, mcrosier, sdardis, nemanjai, nhaehnle, javed.absar, sbc100, jgravelle-google, aheejin, kbarton, llvm-commits Differential Revision: https://reviews.llvm.org/D41464 git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@321234 91177308-0d34-0410-b5e6-96231b3b80d8
2017-12-18[Memcpy Loop Lowering] Remove the fixed int8 lowering.Sean Fertile
Switch over to the lowering that uses target supplied operand types. Differential Revision: https://reviews.llvm.org/D41201 git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@320989 91177308-0d34-0410-b5e6-96231b3b80d8
2017-12-15MachineFunction: Return reference from getFunction(); NFCMatthias Braun
The Function can never be nullptr so we can return a reference. git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@320884 91177308-0d34-0410-b5e6-96231b3b80d8
2017-12-14TLI: Allow using PSV for intrinsic mem operandsMatt Arsenault
git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@320756 91177308-0d34-0410-b5e6-96231b3b80d8
2017-12-14DAG: Expose all MMO flags in getTgtMemIntrinsicMatt Arsenault
Rather than adding more bits to express every MMO flag you could want, just directly use the MMO flags. Also fixes using a bunch of bool arguments to getMemIntrinsicNode. On AMDGPU, buffer and image intrinsics should always have MODereferencable set, but currently there is no way to do that directly during the initial intrinsic lowering. git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@320746 91177308-0d34-0410-b5e6-96231b3b80d8
2017-12-07[CodeGen] Use MachineOperand::print in the MIRPrinter for MO_Register.Francis Visoiu Mistrih
Work towards the unification of MIR and debug output by refactoring the interfaces. For MachineOperand::print, keep a simple version that can be easily called from `dump()`, and a more complex one which will be called from both the MIRPrinter and MachineInstr::print. Add extra checks inside MachineOperand for detached operands (operands with getParent() == nullptr). https://reviews.llvm.org/D40836 * find . \( -name "*.mir" -o -name "*.cpp" -o -name "*.h" -o -name "*.ll" -o -name "*.s" \) -type f -print0 | xargs -0 sed -i '' -E 's/kill: ([^ ]+) ([^ ]+)<def> ([^ ]+)/kill: \1 def \2 \3/g' * find . \( -name "*.mir" -o -name "*.cpp" -o -name "*.h" -o -name "*.ll" -o -name "*.s" \) -type f -print0 | xargs -0 sed -i '' -E 's/kill: ([^ ]+) ([^ ]+) ([^ ]+)<def>/kill: \1 \2 def \3/g' * find . \( -name "*.mir" -o -name "*.cpp" -o -name "*.h" -o -name "*.ll" -o -name "*.s" \) -type f -print0 | xargs -0 sed -i '' -E 's/kill: def ([^ ]+) ([^ ]+) ([^ ]+)<def>/kill: def \1 \2 def \3/g' * find . \( -name "*.mir" -o -name "*.cpp" -o -name "*.h" -o -name "*.ll" -o -name "*.s" \) -type f -print0 | xargs -0 sed -i '' -E 's/<def>//g' * find . \( -name "*.mir" -o -name "*.cpp" -o -name "*.h" -o -name "*.ll" -o -name "*.s" \) -type f -print0 | xargs -0 sed -i '' -E 's/([^ ]+)<kill>/killed \1/g' * find . \( -name "*.mir" -o -name "*.cpp" -o -name "*.h" -o -name "*.ll" -o -name "*.s" \) -type f -print0 | xargs -0 sed -i '' -E 's/([^ ]+)<imp-use,kill>/implicit killed \1/g' * find . \( -name "*.mir" -o -name "*.cpp" -o -name "*.h" -o -name "*.ll" -o -name "*.s" \) -type f -print0 | xargs -0 sed -i '' -E 's/([^ ]+)<dead>/dead \1/g' * find . \( -name "*.mir" -o -name "*.cpp" -o -name "*.h" -o -name "*.ll" -o -name "*.s" \) -type f -print0 | xargs -0 sed -i '' -E 's/([^ ]+)<def[ ]*,[ ]*dead>/dead \1/g' * find . \( -name "*.mir" -o -name "*.cpp" -o -name "*.h" -o -name "*.ll" -o -name "*.s" \) -type f -print0 | xargs -0 sed -i '' -E 's/([^ ]+)<imp-def[ ]*,[ ]*dead>/implicit-def dead \1/g' * find . \( -name "*.mir" -o -name "*.cpp" -o -name "*.h" -o -name "*.ll" -o -name "*.s" \) -type f -print0 | xargs -0 sed -i '' -E 's/([^ ]+)<imp-def>/implicit-def \1/g' * find . \( -name "*.mir" -o -name "*.cpp" -o -name "*.h" -o -name "*.ll" -o -name "*.s" \) -type f -print0 | xargs -0 sed -i '' -E 's/([^ ]+)<imp-use>/implicit \1/g' * find . \( -name "*.mir" -o -name "*.cpp" -o -name "*.h" -o -name "*.ll" -o -name "*.s" \) -type f -print0 | xargs -0 sed -i '' -E 's/([^ ]+)<internal>/internal \1/g' * find . \( -name "*.mir" -o -name "*.cpp" -o -name "*.h" -o -name "*.ll" -o -name "*.s" \) -type f -print0 | xargs -0 sed -i '' -E 's/([^ ]+)<undef>/undef \1/g' git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@320022 91177308-0d34-0410-b5e6-96231b3b80d8
2017-12-06[NVPTX,CUDA] Added llvm.nvvm.fns intrinsic and matching __nvvm_fns builtin ↵Artem Belevich
in clang. Differential Revision: https://reviews.llvm.org/D40872 git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@319909 91177308-0d34-0410-b5e6-96231b3b80d8
2017-12-04[NVPTX] Assign valid global namesJonas Hahnfeld
PTX requires that identifiers consist only of [a-zA-Z0-9_$]. The existing pass already ensured this for globals and this patch adds the cleanup for functions with local linkage. However, there was a different problem in the case of collisions of the adjusted name: The ValueSymbolTable then automatically appended ".N" with increasing Ns to get a unique name while helping the ABI demangling. Special case this behavior to omit the dots and append N directly. This will always give us legal names according to the PTX requirements. Differential Revision: https://reviews.llvm.org/D40573 git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@319657 91177308-0d34-0410-b5e6-96231b3b80d8
2017-11-30[CodeGen] Print "%vreg0" as "%0" in both MIR and debug outputFrancis Visoiu Mistrih
As part of the unification of the debug format and the MIR format, avoid printing "vreg" for virtual registers (which is one of the current MIR possibilities). Basically: * find . \( -name "*.mir" -o -name "*.cpp" -o -name "*.h" -o -name "*.ll" \) -type f -print0 | xargs -0 sed -i '' -E "s/%vreg([0-9]+)/%\1/g" * grep -nr '%vreg' . and fix if needed * find . \( -name "*.mir" -o -name "*.cpp" -o -name "*.h" -o -name "*.ll" \) -type f -print0 | xargs -0 sed -i '' -E "s/ vreg([0-9]+)/ %\1/g" * grep -nr 'vreg[0-9]\+' . and fix if needed Differential Revision: https://reviews.llvm.org/D40420 git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@319427 91177308-0d34-0410-b5e6-96231b3b80d8
2017-11-17Fix a bunch more layering of CodeGen headers that are in TargetDavid Blaikie
All these headers already depend on CodeGen headers so moving them into CodeGen fixes the layering (since CodeGen depends on Target, not the other way around). git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@318490 91177308-0d34-0410-b5e6-96231b3b80d8
2017-11-15Add backend name to Target to enable runtime info to be fed back into TableGenDaniel Sanders
Summary: Make it possible to feed runtime information back to tablegen to enable profile-guided tablegen-eration, detection of untested tablegen definitions, etc. Being a cross-compiler by nature, LLVM will potentially collect data for multiple architectures (e.g. when running 'ninja check'). We therefore need a way for TableGen to figure out what data applies to the backend it is generating at the time. This patch achieves that by including the name of the 'def X : Target ...' for the backend in the TargetRegistry. Reviewers: qcolombet Reviewed By: qcolombet Subscribers: jholewinski, arsenm, jyknight, aditya_nandakumar, sdardis, nemanjai, ab, nhaehnle, t.p.northover, javed.absar, qcolombet, llvm-commits, fedor.sergeev Differential Revision: https://reviews.llvm.org/D39742 git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@318352 91177308-0d34-0410-b5e6-96231b3b80d8
2017-11-14Fix unused variable warning.Richard Smith
git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@318201 91177308-0d34-0410-b5e6-96231b3b80d8
2017-11-14Mark intrinsics operating on the whole warp as IntrInaccessibleMemOnlyArtem Belevich
It's needed to model the fact that they do access data from other threads in a warp and thus can't be CSE'd. git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@318173 91177308-0d34-0410-b5e6-96231b3b80d8
2017-11-08Target/TargetInstrInfo.h -> CodeGen/TargetInstrInfo.h to match layeringDavid Blaikie
This header includes CodeGen headers, and is not, itself, included by any Target headers, so move it into CodeGen to match the layering of its implementation. git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@317647 91177308-0d34-0410-b5e6-96231b3b80d8
2017-11-07[NVPTX] Implement __nvvm_atom_add_gen_d builtin.Justin Lebar
Summary: This just seems to have been an oversight. We already supported the f64 atomic add with an explicit scope (e.g. "cta"), but not the scopeless version. Reviewers: tra Subscribers: jholewinski, sanjoy, cfe-commits, llvm-commits, hiraditya Differential Revision: https://reviews.llvm.org/D39638 git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@317623 91177308-0d34-0410-b5e6-96231b3b80d8
2017-11-03Move TargetFrameLowering.h to CodeGen where it's implementedDavid Blaikie
This header already includes a CodeGen header and is implemented in lib/CodeGen, so move the header there to match. This fixes a link error with modular codegeneration builds - where a header and its implementation are circularly dependent and so need to be in the same library, not split between two like this. git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@317379 91177308-0d34-0410-b5e6-96231b3b80d8
2017-10-24[NVPTX] allow address space inference for volatile loads/stores.Artem Belevich
If particular target supports volatile memory access operations, we can avoid AS casting to generic AS. Currently it's only enabled in NVPTX for loads and stores that access global & shared AS. Differential Revision: https://reviews.llvm.org/D39026 git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@316495 91177308-0d34-0410-b5e6-96231b3b80d8
2017-10-12Revert "TargetMachine: Merge TargetMachine and LLVMTargetMachine"Matthias Braun
Reverting to investigate layering effects of MCJIT not linking libCodeGen but using TargetMachine::getNameWithPrefix() breaking the lldb bots. This reverts commit r315633. git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@315637 91177308-0d34-0410-b5e6-96231b3b80d8
2017-10-12TargetMachine: Merge TargetMachine and LLVMTargetMachineMatthias Braun
Merge LLVMTargetMachine into TargetMachine. - There is no in-tree target anymore that just implements TargetMachine but not LLVMTargetMachine. - It should still be possible to stub out all the various functions in case a target does not want to use lib/CodeGen - This simplifies the code and avoids methods ending up in the wrong interface. Differential Revision: https://reviews.llvm.org/D38489 git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@315633 91177308-0d34-0410-b5e6-96231b3b80d8
2017-10-12[NVPTX] Implemented wmma intrinsics and instructions.Artem Belevich
WMMA = "Warp Level Matrix Multiply-Accumulate". These are the new instructions introduced in PTX6.0 and available on sm_70 GPUs. Differential Revision: https://reviews.llvm.org/D38645 git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@315601 91177308-0d34-0410-b5e6-96231b3b80d8
2017-10-02Remove dead file.Michael Liao
git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@314720 91177308-0d34-0410-b5e6-96231b3b80d8
2017-09-26[NVPTX] added match.{any,all}.sync instructions, intrinsics & builtins.Artem Belevich
Differential Revision: https://reviews.llvm.org/D38191 git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@314223 91177308-0d34-0410-b5e6-96231b3b80d8
2017-09-25Revert "[NVPTX] added match.{any,all}.sync instructions, intrinsics & ↵Justin Lebar
builtins.", rL314135. Causing assertion failures on macos: > Assertion failed: (Num < NumOperands && "Invalid child # of SDNode!"), > function getOperand, file > /Users/buildslave/jenkins/workspace/clang-stage1-cmake-RA-incremental/llvm/include/llvm/CodeGen/SelectionDAGNodes.h, > line 835. http://green.lab.llvm.org/green/job/clang-stage1-cmake-RA-incremental/42739/testReport/LLVM/CodeGen_NVPTX/surf_read_cuda_ll/ git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@314142 91177308-0d34-0410-b5e6-96231b3b80d8
2017-09-25[NVPTX] added match.{any,all}.sync instructions, intrinsics & builtins.Artem Belevich
Differential Revision: https://reviews.llvm.org/D38191 git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@314135 91177308-0d34-0410-b5e6-96231b3b80d8
2017-09-21[NVPTX] Implemented bar.warp.sync, barrier.sync, and vote{.sync} ↵Artem Belevich
instructions/intrinsics/builtins. Differential Revision: https://reviews.llvm.org/D38148 git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@313898 91177308-0d34-0410-b5e6-96231b3b80d8
2017-09-20[NVPTX] Implemented shfl.sync instruction and supporting intrinsics/builtins.Artem Belevich
Differential Revision: https://reviews.llvm.org/D38090 git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@313820 91177308-0d34-0410-b5e6-96231b3b80d8
2017-09-07[CUDA] Added rudimentary support for CUDA-9 and sm_70.Artem Belevich
For now CUDA-9 is not included in the list of CUDA versions clang searches for, so the path to CUDA-9 must be explicitly passed via --cuda-path=. On LLVM side NVPTX added sm_70 GPU type which bumps required PTX version to 6.0, but otherwise is equivalent to sm_62 at the moment. Differential Revision: https://reviews.llvm.org/D37576 git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@312734 91177308-0d34-0410-b5e6-96231b3b80d8
2017-08-20[NVPTX] Reduce copypasta.Benjamin Kramer
No functionality change intended. git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@311295 91177308-0d34-0410-b5e6-96231b3b80d8
2017-08-03Delete Default and JITDefault code modelsRafael Espindola
IMHO it is an antipattern to have a enum value that is Default. At any given piece of code it is not clear if we have to handle Default or if has already been mapped to a concrete value. In this case in particular, only the target can do the mapping and it is nice to make sure it is always done. This deletes the two default enum values of CodeModel and uses an explicit Optional<CodeModel> when it is possible that it is unspecified. git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@309911 91177308-0d34-0410-b5e6-96231b3b80d8
2017-07-26Change CallLoweringInfo::CS to be an ImmutableCallSite instead of a pointer. ↵Peter Collingbourne
NFCI. This was a use-after-free waiting to happen. git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@309159 91177308-0d34-0410-b5e6-96231b3b80d8
2017-07-21[SystemZ, LoopStrengthReduce]Jonas Paulsson
This patch makes LSR generate better code for SystemZ in the cases of memory intrinsics, Load->Store pairs or comparison of immediate with memory. In order to achieve this, the following common code changes were made: * New TTI hook: LSRWithInstrQueries(), which defaults to false. Controls if LSR should do instruction-based addressing evaluations by calling isLegalAddressingMode() with the Instruction pointers. * In LoopStrengthReduce: handle address operands of memset, memmove and memcpy as address uses, and call isFoldableMemAccessOffset() for any LSRUse::Address, not just loads or stores. SystemZ changes: * isLSRCostLess() implemented with Insns first, and without ImmCost. * New function supportedAddressingMode() that is a helper for TTI methods looking at Instructions passed via pointers. Review: Ulrich Weigand, Quentin Colombet https://reviews.llvm.org/D35262 https://reviews.llvm.org/D35049 git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@308729 91177308-0d34-0410-b5e6-96231b3b80d8
2017-07-20[NVPTX] Add lowering of i128 params.Artem Belevich
The patch adds support of i128 params lowering. The changes are quite trivial to support i128 as a "special case" of integer type. With this patch, we lower i128 params the same way as aggregates of size 16 bytes: .param .b8 _ [16]. Currently, NVPTX can't deal with the 128 bit integers: * in some cases because of failed assertions like ValVTs.size() == OutVals.size() && "Bad return value decomposition" * in other cases emitting PTX with .i128 or .u128 types (which are not valid [1]) [1] http://docs.nvidia.com/cuda/parallel-thread-execution/index.html#fundamental-types Differential Revision: https://reviews.llvm.org/D34555 Patch by: Denys Zariaiev (denys.zariaiev@gmail.com) git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@308675 91177308-0d34-0410-b5e6-96231b3b80d8
2017-07-20Changed EOL back to LF. NFC.Artem Belevich
git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@308671 91177308-0d34-0410-b5e6-96231b3b80d8
2017-07-12Remove unneeded use of #undef DEBUG_TYPE. NFCSam Clegg
Where is is needed (at the end of headers that define it), be consistent about its use. Also fix a few header guards that I found in the process. Differential Revision: https://reviews.llvm.org/D34916 git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@307840 91177308-0d34-0410-b5e6-96231b3b80d8
2017-07-07Extend memcpy expansion in Transform/Utils to handle wider operand types.Sean Fertile
Adds loop expansions for known-size and unknown-sized memcpy calls, allowing the target to provide the operand types through TTI callbacks. The default values for the TTI callbacks use int8 operand types and matches the existing behaviour if they aren't overridden by the target. Differential revision: https://reviews.llvm.org/D32536 git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@307346 91177308-0d34-0410-b5e6-96231b3b80d8
2017-07-06Reverting r307326 because it breaks clang tests.Michael Kuperstein
git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@307334 91177308-0d34-0410-b5e6-96231b3b80d8
2017-07-06[NVPTX] Add lowering of i128 params.Michael Kuperstein
The patch adds support of i128 params lowering. The changes are quite trivial to support i128 as a "special case" of integer type. With this patch, we lower i128 params the same way as aggregates of size 16 bytes: .param .b8 _ [16]. Currently, NVPTX can't deal with the 128 bit integers: * in some cases because of failed assertions like ValVTs.size() == OutVals.size() && "Bad return value decomposition" * in other cases emitting PTX with .i128 or .u128 types (which are not valid [1]) [1] http://docs.nvidia.com/cuda/parallel-thread-execution/index.html#fundamental-types Differential Revision: https://reviews.llvm.org/D34555 Patch by: Denys Zariaiev (denys.zariaiev@gmail.com) git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@307326 91177308-0d34-0410-b5e6-96231b3b80d8
2017-07-04fix trivial typos in comments; NFCHiroshi Inoue
git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@307094 91177308-0d34-0410-b5e6-96231b3b80d8
2017-06-28[LoopUnroll] Pass SCEV to getUnrollingPreferences hook. NFCI.Geoff Berry
Reviewers: sanjoy, anna, reames, apilipenko, igor-laevsky, mkuper Subscribers: jholewinski, arsenm, mzolotukhin, nemanjai, nhaehnle, javed.absar, mcrosier, llvm-commits Differential Revision: https://reviews.llvm.org/D34531 git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@306554 91177308-0d34-0410-b5e6-96231b3b80d8
2017-06-06Sort the remaining #include lines in include/... and lib/....Chandler Carruth
I did this a long time ago with a janky python script, but now clang-format has built-in support for this. I fed clang-format every line with a #include and let it re-sort things according to the precise LLVM rules for include ordering baked into clang-format these days. I've reverted a number of files where the results of sorting includes isn't healthy. Either places where we have legacy code relying on particular include ordering (where possible, I'll fix these separately) or where we have particular formatting around #include lines that I didn't want to disturb in this patch. This patch is *entirely* mechanical. If you get merge conflicts or anything, just ignore the changes in this patch and run clang-format over your #include lines in the files. Sorry for any noise here, but it is important to keep these things stable. I was seeing an increasing number of patches with irrelevant re-ordering of #include lines because clang-format was used. This patch at least isolates that churn, makes it easy to skip when resolving conflicts, and gets us to a clean baseline (again). git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@304787 91177308-0d34-0410-b5e6-96231b3b80d8
2017-05-31TargetMachine: Indicate whether machine verifier passes.Matthias Braun
This adds a callback to the LLVMTargetMachine that lets target indicate that they do not pass the machine verifier checks in all cases yet. This is intended to be a temporary measure while the targets are fixed allowing us to enable the machine verifier by default with EXPENSIVE_CHECKS enabled! Differential Revision: https://reviews.llvm.org/D33696 git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@304320 91177308-0d34-0410-b5e6-96231b3b80d8
2017-05-30TargetPassConfig: Keep a reference to an LLVMTargetMachine; NFCMatthias Braun
TargetPassConfig is not useful for targets that do not use the CodeGen library, so we may just as well store a pointer to an LLVMTargetMachine instead of just to a TargetMachine. While at it, also change the constructor to take a reference instead of a pointer as the TM must not be nullptr. git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@304247 91177308-0d34-0410-b5e6-96231b3b80d8
2017-05-15[NVPTX] Don't flag StoreParam/LoadParam memory chain operands as ↵Simon Pilgrim
ReadMem/WriteMem (PR32146) Follow up to D33147 NVPTXTargetLowering::LowerCall was trusting the default argument values. Fixes another 17 of the NVPTX '-verify-machineinstrs with EXPENSIVE_CHECKS' errors in PR32146. Differential Revision: https://reviews.llvm.org/D33189 git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@303082 91177308-0d34-0410-b5e6-96231b3b80d8
2017-05-15[NVPTX] Don't rely on default arguments to ↵Simon Pilgrim
SelectionDAG::getMemIntrinsicNode. NFC. NFC followup to D33147, this explicitly sets all the arguments (instead of relying on the defaults) to SelectionDAG::getMemIntrinsicNode to help identify -verify-machineinstrs issues. git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@303047 91177308-0d34-0410-b5e6-96231b3b80d8
2017-05-12[NVPTX] Don't flag StoreRetVal memory chain operands as ReadMem (PR32146)Simon Pilgrim
This fixes 47 of the 75 NVPTX '-verify-machineinstrs with EXPENSIVE_CHECKS' errors in PR32146. Differential Revision: https://reviews.llvm.org/D33147 git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@302942 91177308-0d34-0410-b5e6-96231b3b80d8
2017-05-09Add extra operand to CALLSEQ_START to keep frame part set up previouslySerge Pavlov
Using arguments with attribute inalloca creates problems for verification of machine representation. This attribute instructs the backend that the argument is prepared in stack prior to CALLSEQ_START..CALLSEQ_END sequence (see http://llvm.org/docs/InAlloca.htm for details). Frame size stored in CALLSEQ_START in this case does not count the size of this argument. However CALLSEQ_END still keeps total frame size, as caller can be responsible for cleanup of entire frame. So CALLSEQ_START and CALLSEQ_END keep different frame size and the difference is treated by MachineVerifier as stack error. Currently there is no way to distinguish this case from actual errors. This patch adds additional argument to CALLSEQ_START and its target-specific counterparts to keep size of stack that is set up prior to the call frame sequence. This argument allows MachineVerifier to calculate actual frame size associated with frame setup instruction and correctly process the case of inalloca arguments. The changes made by the patch are: - Frame setup instructions get the second mandatory argument. It affects all targets that use frame pseudo instructions and touched many files although the changes are uniform. - Access to frame properties are implemented using special instructions rather than calls getOperand(N).getImm(). For X86 and ARM such replacement was made previously. - Changes that reflect appearance of additional argument of frame setup instruction. These involve proper instruction initialization and methods that access instruction arguments. - MachineVerifier retrieves frame size using method, which reports sum of frame parts initialized inside frame instruction pair and outside it. The patch implements approach proposed by Quentin Colombet in https://bugs.llvm.org/show_bug.cgi?id=27481#c1. It fixes 9 tests failed with machine verifier enabled and listed in PR27481. Differential Revision: https://reviews.llvm.org/D32394 git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@302527 91177308-0d34-0410-b5e6-96231b3b80d8
2017-05-06[NVPTX] Add support for ISD::ABS loweringSimon Pilgrim
Use the ISD::ABS opcode directly Differential Revision: https://reviews.llvm.org/D32944 git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@302356 91177308-0d34-0410-b5e6-96231b3b80d8
generated by cgit v1.2.3 (git 2.25.1) at 2024-06-05 16:22:05 +0000