summaryrefslogtreecommitdiffstats
path: root/include/clang/Basic/BuiltinsNVPTX.def
Commit message (Collapse)AuthorAgeFilesLines
* [CUDA] Removed unused __nvvm_* builtins with non-generic pointers.Artem Belevich2018-06-201-74/+0
| | | | | | | | | | | | They were hot even hooked into CGBuiltin's machinery. Even if they were, CUDA does not support AS-specific pointers, so there would be no legal way no way to call these builtins. This came up in D47154. Differential Revision: https://reviews.llvm.org/D47845 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@335168 91177308-0d34-0410-b5e6-96231b3b80d8
* [NVPTX, CUDA] Added support for m8n32k16 and m32n8k16 variants of wmma ↵Artem Belevich2018-04-181-13/+47
| | | | | | | | | | instructions. The new instructions were added added for sm_70+ GPUs in CUDA-9.1. Differential Revision: https://reviews.llvm.org/D45068 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@330296 91177308-0d34-0410-b5e6-96231b3b80d8
* [NVPTX, CUDA] Improved feature constraints on NVPTX target builtins.Artem Belevich2018-04-111-100/+108
| | | | | | | | | | 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,CUDA] Added llvm.nvvm.fns intrinsic and matching __nvvm_fns builtin ↵Artem Belevich2017-12-061-0/+3
| | | | | | | | in clang. Differential Revision: https://reviews.llvm.org/D40872 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@319909 91177308-0d34-0410-b5e6-96231b3b80d8
* [NVPTX] Implement __nvvm_atom_add_gen_d builtin.Justin Lebar2017-11-071-1/+1
| | | | | | | | | | | | | | | 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/cfe/trunk@317623 91177308-0d34-0410-b5e6-96231b3b80d8
* [CUDA] Added __hmma_m16n16k16_* builtins to support mma instructions on sm_70Artem Belevich2017-10-121-0/+13
| | | | | | Differential Revision: https://reviews.llvm.org/D38742 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@315624 91177308-0d34-0410-b5e6-96231b3b80d8
* [NVPTX] added match.{any,all}.sync instructions, intrinsics & builtins.Artem Belevich2017-09-261-0/+7
| | | | | | Differential Revision: https://reviews.llvm.org/D38191 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@314223 91177308-0d34-0410-b5e6-96231b3b80d8
* Revert "[NVPTX] added match.{any,all}.sync instructions, intrinsics & ↵Justin Lebar2017-09-251-7/+0
| | | | | | | | | | | | | | | 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/cfe/trunk@314142 91177308-0d34-0410-b5e6-96231b3b80d8
* [NVPTX] added match.{any,all}.sync instructions, intrinsics & builtins.Artem Belevich2017-09-251-0/+7
| | | | | | Differential Revision: https://reviews.llvm.org/D38191 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@314135 91177308-0d34-0410-b5e6-96231b3b80d8
* [NVPTX] Implemented bar.warp.sync, barrier.sync, and vote{.sync} ↵Artem Belevich2017-09-211-0/+14
| | | | | | | | 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/+9
| | | | | | Differential Revision: https://reviews.llvm.org/D38090 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@313820 91177308-0d34-0410-b5e6-96231b3b80d8
* [NVPTX] Auto-upgrade some NVPTX intrinsics to LLVM target-generic code.Justin Lebar2017-01-211-24/+0
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | Summary: Specifically, we upgrade llvm.nvvm.: * brev{32,64} * clz.{i,ll} * popc.{i,ll} * abs.{i,ll} * {min,max}.{i,ll,u,ull} * h2f These either map directly to an existing LLVM target-generic intrinsic or map to a simple LLVM target-generic idiom. In all cases, we check that the code we generate is lowered to PTX as we expect. These builtins don't need to be backfilled in clang: They're not accessible to user code from nvcc. Reviewers: tra Subscribers: majnemer, cfe-commits, llvm-commits, jholewinski Differential Revision: https://reviews.llvm.org/D28793 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@292694 91177308-0d34-0410-b5e6-96231b3b80d8
* [CUDA] added __nvvm_atom_{sys|cta}_* builtins.Artem Belevich2016-09-281-0/+73
| | | | | | | | 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-47/+44
| | | | | | 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
* NVPTX: Rename __builtin_ptx_shfl -> __nvvm_shflJustin Bogner2016-07-061-8/+8
| | | | | | | To match "NVPTX: Make the llvm.nvvm.shfl intrinsics and builtin names consistent" in LLVM. git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@274663 91177308-0d34-0410-b5e6-96231b3b80d8
* [CUDA] Implement __shfl* intrinsics in clang headers.Justin Lebar2016-06-091-0/+11
| | | | | | | | | | | | Summary: Clang changes to make use of the LLVM intrinsics added in D21160. Reviewers: tra Subscribers: jholewinski, cfe-commits Differential Revision: http://reviews.llvm.org/D21162 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@272299 91177308-0d34-0410-b5e6-96231b3b80d8
* [CUDA] Implement __ldg using intrinsics.Justin Lebar2016-05-191-0/+36
| | | | | | | | | | | | | | | | | | 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] 32-bit NVPTX should have 32-bit long type.Artem Belevich2015-09-281-1/+1
| | | | | | | | | 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] Added stubs for __nvvm_atom_add_*_d() builtins.Artem Belevich2015-08-191-0/+3
| | | | | | | | | They show up in CUDA headers but are not currently supported by NVPTX back-end. Differential Revision: http://reviews.llvm.org/D11694 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@245502 91177308-0d34-0410-b5e6-96231b3b80d8
* [NVPTX] Fixed a typo in __nvvm_atom_min_gen_l() type string.Artem Belevich2015-06-251-1/+1
| | | | | | Differential Revision: http://reviews.llvm.org/D10664 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@240659 91177308-0d34-0410-b5e6-96231b3b80d8
* [NVPTX] Fix type error for some builtins in BuiltinsNVPTX.defJustin Holewinski2014-12-021-15/+15
| | | | git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@223116 91177308-0d34-0410-b5e6-96231b3b80d8
* [NVPTX] Add entire list of supported builtinsJustin Holewinski2013-05-221-243/+501
| | | | git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@182468 91177308-0d34-0410-b5e6-96231b3b80d8
* [NVPTX] Fix const modifier on builtinsJustin Holewinski2012-11-121-7/+7
| | | | | | Some NVVM intrinsics were incorrectly labeled. git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@167700 91177308-0d34-0410-b5e6-96231b3b80d8
* [NVPTX] Add __nvvm_* intrinsics as Clang builtinsJustin Holewinski2012-11-091-0/+246
| | | | | | Fixes bug 13354. git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@167647 91177308-0d34-0410-b5e6-96231b3b80d8
* Replace PTX back-end with NVPTX back-end in all places where Clang caresJustin Holewinski2012-05-241-0/+62
NV_CONTRIB git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@157403 91177308-0d34-0410-b5e6-96231b3b80d8
generated by cgit v1.2.3 (git 2.25.1) at 2024-05-03 21:29:32 +0000