diff options
| author | Justin Lebar <jlebar@google.com> | 2016-03-22 00:09:28 +0000 |
|---|---|---|
| committer | Justin Lebar <jlebar@google.com> | 2016-03-22 00:09:28 +0000 |
| commit | 717d2b0a0dd420fa7734396b89b22c03bb572be9 (patch) | |
| tree | 442a094e96ebdd14b5554a9719d14bbf49837cb8 /clang/test/CodeGen | |
| parent | e6a2cc12a33835b9e3889c6bbb46a62065420f89 (diff) | |
| download | bcm5719-llvm-717d2b0a0dd420fa7734396b89b22c03bb572be9.tar.gz bcm5719-llvm-717d2b0a0dd420fa7734396b89b22c03bb572be9.zip | |
[CUDA] Implement atomicInc and atomicDec builtins
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
llvm-svn: 264009
Diffstat (limited to 'clang/test/CodeGen')
| -rw-r--r-- | clang/test/CodeGen/builtins-nvptx.c | 8 |
1 files changed, 7 insertions, 1 deletions
diff --git a/clang/test/CodeGen/builtins-nvptx.c b/clang/test/CodeGen/builtins-nvptx.c index 745e74f0ca6..1a6f889327e 100644 --- a/clang/test/CodeGen/builtins-nvptx.c +++ b/clang/test/CodeGen/builtins-nvptx.c @@ -189,7 +189,7 @@ __shared__ long long sll; // Check for atomic intrinsics // CHECK-LABEL: nvvm_atom -__device__ void nvvm_atom(float *fp, float f, int *ip, int i, long *lp, long l, +__device__ void nvvm_atom(float *fp, float f, int *ip, int i, unsigned int *uip, unsigned ui, long *lp, long l, long long *llp, long long ll) { // CHECK: atomicrmw add __nvvm_atom_add_gen_i(ip, i); @@ -272,5 +272,11 @@ __device__ void nvvm_atom(float *fp, float f, int *ip, int i, long *lp, long l, // CHECK: call float @llvm.nvvm.atomic.load.add.f32.p0f32 __nvvm_atom_add_gen_f(fp, f); + // CHECK: call i32 @llvm.nvvm.atomic.load.inc.32.p0i32 + __nvvm_atom_inc_gen_ui(uip, ui); + + // CHECK: call i32 @llvm.nvvm.atomic.load.dec.32.p0i32 + __nvvm_atom_dec_gen_ui(uip, ui); + // CHECK: ret } |
