summaryrefslogtreecommitdiff
path: root/test/CodeGen/builtins-nvptx.c
Commit message (Collapse)AuthorAgeFilesLines
* [CodeGen] NVPTX: Switch from atomic.load.add.f32 to atomicrmw faddBenjamin Kramer2019-07-111-1/+1
| | | | git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@365798 91177308-0d34-0410-b5e6-96231b3b80d8
* [NVPTX, CUDA] Improved feature constraints on NVPTX target builtins.Artem Belevich2018-04-111-68/+71
| | | | | | | | | | When NVPTX TARGET_BUILTIN specifies sm_XX or ptxYY as required feature, consider those features available if we're compiling for GPU >= sm_XX or have enabled PTX version >= ptxYY. Differential Revision: https://reviews.llvm.org/D45061 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@329829 91177308-0d34-0410-b5e6-96231b3b80d8
* [NVPTX] Implemented bar.warp.sync, barrier.sync, and vote{.sync} ↵Artem Belevich2017-09-211-0/+12
| | | | | | | | instructions/intrinsics/builtins. Differential Revision: https://reviews.llvm.org/D38148 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@313898 91177308-0d34-0410-b5e6-96231b3b80d8
* [NVPTX] Implemented shfl.sync instruction and supporting intrinsics/builtins.Artem Belevich2017-09-201-0/+21
| | | | | | Differential Revision: https://reviews.llvm.org/D38090 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@313820 91177308-0d34-0410-b5e6-96231b3b80d8
* [CUDA] added __nvvm_atom_{sys|cta}_* builtins.Artem Belevich2016-09-281-6/+260
| | | | | | | | These builtins are available on sm_60+ GPU only. Differential Revision: https://reviews.llvm.org/D24944 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@282609 91177308-0d34-0410-b5e6-96231b3b80d8
* [CUDA] Rename the __nvvm_bar0 builtin back to __syncthreads.Justin Lebar2016-07-071-1/+1
| | | | | | | | | The builtin was renamed in r274770. But __syncthreads is part of our user-facing API, so we need to keep the name as-is. Patch by Justin Bogner. git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@274780 91177308-0d34-0410-b5e6-96231b3b80d8
* NVPTX: Use the nvvm builtins to read SRegs rather than the legacy ptx onesJustin Bogner2016-07-071-68/+68
| | | | | | The ptx spellings were removed from LLVM in r274769. git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@274770 91177308-0d34-0410-b5e6-96231b3b80d8
* [CUDA] Implement __ldg using intrinsics.Justin Lebar2016-05-191-2/+104
| | | | | | | | | | | | | | | | | | Summary: Previously it was implemented as inline asm in the CUDA headers. This change allows us to use the [addr+imm] addressing mode when executing ld.global.nc instructions. This translates into a 1.3x speedup on some benchmarks that call this instruction from within an unrolled loop. Reviewers: tra, rsmith Subscribers: jhen, cfe-commits, jholewinski Differential Revision: http://reviews.llvm.org/D19990 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@270150 91177308-0d34-0410-b5e6-96231b3b80d8
* [CUDA] Implement atomicInc and atomicDec builtinsJustin Lebar2016-03-221-1/+7
| | | | | | | | | | | | | | | These functions cannot be implemented as atomicrmw or cmpxchg instructions, so they are implemented as a call to the NVVM intrinsics @llvm.nvvm.atomic.load.inc.32.p0i32 and @llvm.nvvm.atomic.load.dec.32.p0i32. Patch by Jason Henline. Reviewers: jlebar Differential Revision: http://reviews.llvm.org/D18322 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@264009 91177308-0d34-0410-b5e6-96231b3b80d8
* [CUDA] fix codegen for __nvvm_atom_cas_*Jingyue Wu2015-09-301-0/+3
| | | | | | | | | | | | Summary: __nvvm_atom_cas_* returns the old value instead of whether the swap succeeds. Reviewers: eliben, tra Subscribers: jholewinski, llvm-commits Differential Revision: http://reviews.llvm.org/D13306 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@248951 91177308-0d34-0410-b5e6-96231b3b80d8
* [CUDA] 32-bit NVPTX should have 32-bit long type.Artem Belevich2015-09-281-4/+3
| | | | | | | | | Currently it's 64-bit which will lead to mismatch between host and device code if we compile for i386. Differential Revision: http://reviews.llvm.org/D13181 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@248753 91177308-0d34-0410-b5e6-96231b3b80d8
* [CUDA] fix codegen for __nvvm_atom_min/max_gen_u*Jingyue Wu2015-08-311-10/+10
| | | | | | | | | | | | Summary: Clang should emit "atomicrmw umin/umax" instead of "atomicrmw min/max". Reviewers: eliben, tra Subscribers: llvm-commits Differential Revision: http://reviews.llvm.org/D12487 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@246455 91177308-0d34-0410-b5e6-96231b3b80d8
* [CUDA] Implemented __nvvm_atom_*_gen_* builtins.Artem Belevich2015-06-251-13/+109
| | | | | | | | | | | Integer variants are implemented as atomicrmw or cmpxchg instructions. Atomic add for floating point (__nvvm_atom_add_gen_f()) is implemented as a call to an overloaded @llvm.nvvm.atomic.load.add.f32.* LVVM intrinsic. Differential Revision: http://reviews.llvm.org/D10666 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@240669 91177308-0d34-0410-b5e6-96231b3b80d8
* [NVPTX] Fix type error for some builtins in BuiltinsNVPTX.defJustin Holewinski2014-12-021-0/+2
| | | | git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@223116 91177308-0d34-0410-b5e6-96231b3b80d8
* clang/test/CodeGen/builtins-nvptx.c: Prune "REQUIRES: ↵NAKAMURA Takumi2013-12-041-1/+0
| | | | | | nvptx64-registered-target". "nvptx" should imply it. git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@196348 91177308-0d34-0410-b5e6-96231b3b80d8
* [NVPTX] Add entire list of supported builtinsJustin Holewinski2013-05-221-0/+9
| | | | git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@182468 91177308-0d34-0410-b5e6-96231b3b80d8
* [NVPTX] Add __nvvm_* intrinsics as Clang builtinsJustin Holewinski2012-11-091-2/+72
| | | | | | Fixes bug 13354. git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@167647 91177308-0d34-0410-b5e6-96231b3b80d8
* Revert r166541, "clang/test: Add appropriate requirements as REQUIRES, ↵NAKAMURA Takumi2012-10-241-1/+0
| | | | | | | | corresponding to r166532." According to r166543, it is not needed for now. git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@166544 91177308-0d34-0410-b5e6-96231b3b80d8
* clang/test: Add appropriate requirements as REQUIRES, corresponding to r166532.NAKAMURA Takumi2012-10-241-0/+1
| | | | git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@166541 91177308-0d34-0410-b5e6-96231b3b80d8
* Un-XFAIL CodeGen/builtins-nvptx.c now that the proper changes haveJustin Holewinski2012-05-241-2/+0
| | | | | | landed in LLVM core git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@157418 91177308-0d34-0410-b5e6-96231b3b80d8
* XFAIL this test, which does not pass on trunk since the grandJohn McCall2012-05-241-0/+1
| | | | | | renaming in r157403. git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@157413 91177308-0d34-0410-b5e6-96231b3b80d8
* Replace PTX back-end with NVPTX back-end in all places where Clang caresJustin Holewinski2012-05-241-0/+99
NV_CONTRIB git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@157403 91177308-0d34-0410-b5e6-96231b3b80d8
View Lorry Controller Status