diff options
Diffstat (limited to 'clang')
| -rw-r--r-- | clang/lib/CodeGen/CGBuiltin.cpp | 4 | ||||
| -rw-r--r-- | clang/test/CodeGen/builtins-nvptx.c | 3 |
2 files changed, 6 insertions, 1 deletions
diff --git a/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp index f5e3aca103a..27dd5b5728d 100644 --- a/clang/lib/CodeGen/CGBuiltin.cpp +++ b/clang/lib/CodeGen/CGBuiltin.cpp @@ -7021,7 +7021,9 @@ Value *CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID, case NVPTX::BI__nvvm_atom_cas_gen_i: case NVPTX::BI__nvvm_atom_cas_gen_l: case NVPTX::BI__nvvm_atom_cas_gen_ll: - return MakeAtomicCmpXchgValue(*this, E, true); + // __nvvm_atom_cas_gen_* should return the old value rather than the + // success flag. + return MakeAtomicCmpXchgValue(*this, E, /*ReturnBool=*/false); case NVPTX::BI__nvvm_atom_add_gen_f: { Value *Ptr = EmitScalarExpr(E->getArg(0)); 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 |
