diff options
| author | Jingyue Wu <jingyue@google.com> | 2015-09-30 21:49:32 +0000 |
|---|---|---|
| committer | Jingyue Wu <jingyue@google.com> | 2015-09-30 21:49:32 +0000 |
| commit | f1eca25b16b9016b2c058559106ec233ceabfa5c (patch) | |
| tree | 97b7e6d1e076e1fe3dcbe306daf6de146dfb2817 /clang/test/CodeGen/builtins-nvptx.c | |
| parent | e5e83474966a67e2fb1964fd847f975a65bc3699 (diff) | |
| download | bcm5719-llvm-f1eca25b16b9016b2c058559106ec233ceabfa5c.tar.gz bcm5719-llvm-f1eca25b16b9016b2c058559106ec233ceabfa5c.zip | |
[CUDA] fix codegen for __nvvm_atom_cas_*
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
llvm-svn: 248951
Diffstat (limited to 'clang/test/CodeGen/builtins-nvptx.c')
| -rw-r--r-- | clang/test/CodeGen/builtins-nvptx.c | 3 |
1 files changed, 3 insertions, 0 deletions
diff --git a/clang/test/CodeGen/builtins-nvptx.c b/clang/test/CodeGen/builtins-nvptx.c index d1abb9c9509..745e74f0ca6 100644 --- a/clang/test/CodeGen/builtins-nvptx.c +++ b/clang/test/CodeGen/builtins-nvptx.c @@ -260,10 +260,13 @@ __device__ void nvvm_atom(float *fp, float f, int *ip, int i, long *lp, long l, __nvvm_atom_min_gen_ull((unsigned long long *)&sll, ll); // CHECK: cmpxchg + // CHECK-NEXT: extractvalue { i32, i1 } {{%[0-9]+}}, 0 __nvvm_atom_cas_gen_i(ip, 0, i); // CHECK: cmpxchg + // CHECK-NEXT: extractvalue { {{i32|i64}}, i1 } {{%[0-9]+}}, 0 __nvvm_atom_cas_gen_l(&dl, 0, l); // CHECK: cmpxchg + // CHECK-NEXT: extractvalue { i64, i1 } {{%[0-9]+}}, 0 __nvvm_atom_cas_gen_ll(&sll, 0, ll); // CHECK: call float @llvm.nvvm.atomic.load.add.f32.p0f32 |
