diff options
Diffstat (limited to 'clang')
| -rw-r--r-- | clang/lib/CodeGen/CGBuiltin.cpp | 8 | ||||
| -rw-r--r-- | clang/test/CodeGen/builtins-nvptx.c | 20 |
2 files changed, 16 insertions, 12 deletions
diff --git a/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp index 955dcfc96ca..6966b03ef3b 100644 --- a/clang/lib/CodeGen/CGBuiltin.cpp +++ b/clang/lib/CodeGen/CGBuiltin.cpp @@ -6985,18 +6985,22 @@ Value *CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID, case NVPTX::BI__nvvm_atom_max_gen_i: case NVPTX::BI__nvvm_atom_max_gen_l: case NVPTX::BI__nvvm_atom_max_gen_ll: + return MakeBinaryAtomicValue(*this, llvm::AtomicRMWInst::Max, E); + case NVPTX::BI__nvvm_atom_max_gen_ui: case NVPTX::BI__nvvm_atom_max_gen_ul: case NVPTX::BI__nvvm_atom_max_gen_ull: - return MakeBinaryAtomicValue(*this, llvm::AtomicRMWInst::Max, E); + return MakeBinaryAtomicValue(*this, llvm::AtomicRMWInst::UMax, E); case NVPTX::BI__nvvm_atom_min_gen_i: case NVPTX::BI__nvvm_atom_min_gen_l: case NVPTX::BI__nvvm_atom_min_gen_ll: + return MakeBinaryAtomicValue(*this, llvm::AtomicRMWInst::Min, E); + case NVPTX::BI__nvvm_atom_min_gen_ui: case NVPTX::BI__nvvm_atom_min_gen_ul: case NVPTX::BI__nvvm_atom_min_gen_ull: - return MakeBinaryAtomicValue(*this, llvm::AtomicRMWInst::Min, E); + return MakeBinaryAtomicValue(*this, llvm::AtomicRMWInst::UMin, E); case NVPTX::BI__nvvm_atom_cas_gen_i: case NVPTX::BI__nvvm_atom_cas_gen_l: diff --git a/clang/test/CodeGen/builtins-nvptx.c b/clang/test/CodeGen/builtins-nvptx.c index ebf20673ddb..498da476d10 100644 --- a/clang/test/CodeGen/builtins-nvptx.c +++ b/clang/test/CodeGen/builtins-nvptx.c @@ -234,30 +234,30 @@ __device__ void nvvm_atom(float *fp, float f, int *ip, int i, long *lp, long l, // CHECK: atomicrmw xchg __nvvm_atom_xchg_gen_ll(&sll, ll); - // CHECK: atomicrmw max + // CHECK: atomicrmw max i32* __nvvm_atom_max_gen_i(ip, i); - // CHECK: atomicrmw max + // CHECK: atomicrmw umax i32* __nvvm_atom_max_gen_ui((unsigned int *)ip, i); // CHECK: atomicrmw max __nvvm_atom_max_gen_l(&dl, l); - // CHECK: atomicrmw max + // CHECK: atomicrmw umax __nvvm_atom_max_gen_ul((unsigned long *)&dl, l); - // CHECK: atomicrmw max + // CHECK: atomicrmw max i64* __nvvm_atom_max_gen_ll(&sll, ll); - // CHECK: atomicrmw max + // CHECK: atomicrmw umax i64* __nvvm_atom_max_gen_ull((unsigned long long *)&sll, ll); - // CHECK: atomicrmw min + // CHECK: atomicrmw min i32* __nvvm_atom_min_gen_i(ip, i); - // CHECK: atomicrmw min + // CHECK: atomicrmw umin i32* __nvvm_atom_min_gen_ui((unsigned int *)ip, i); // CHECK: atomicrmw min __nvvm_atom_min_gen_l(&dl, l); - // CHECK: atomicrmw min + // CHECK: atomicrmw umin __nvvm_atom_min_gen_ul((unsigned long *)&dl, l); - // CHECK: atomicrmw min + // CHECK: atomicrmw min i64* __nvvm_atom_min_gen_ll(&sll, ll); - // CHECK: atomicrmw min + // CHECK: atomicrmw umin i64* __nvvm_atom_min_gen_ull((unsigned long long *)&sll, ll); // CHECK: cmpxchg |
