Author: jingyue Date: Mon Aug 31 12:25:51 2015 New Revision: 246455 URL: http://llvm.org/viewvc/llvm-project?rev=246455&view=rev Log: [CUDA] fix codegen for __nvvm_atom_min/max_gen_u*
Summary: Clang should emit "atomicrmw umin/umax" instead of "atomicrmw min/max". Reviewers: eliben, tra Subscribers: llvm-commits Differential Revision: http://reviews.llvm.org/D12487 Modified: cfe/trunk/lib/CodeGen/CGBuiltin.cpp cfe/trunk/test/CodeGen/builtins-nvptx.c Modified: cfe/trunk/lib/CodeGen/CGBuiltin.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGBuiltin.cpp?rev=246455&r1=246454&r2=246455&view=diff ============================================================================== --- cfe/trunk/lib/CodeGen/CGBuiltin.cpp (original) +++ cfe/trunk/lib/CodeGen/CGBuiltin.cpp Mon Aug 31 12:25:51 2015 @@ -6985,18 +6985,22 @@ Value *CodeGenFunction::EmitNVPTXBuiltin 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: Modified: cfe/trunk/test/CodeGen/builtins-nvptx.c URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGen/builtins-nvptx.c?rev=246455&r1=246454&r2=246455&view=diff ============================================================================== --- cfe/trunk/test/CodeGen/builtins-nvptx.c (original) +++ cfe/trunk/test/CodeGen/builtins-nvptx.c Mon Aug 31 12:25:51 2015 @@ -234,30 +234,30 @@ __device__ void nvvm_atom(float *fp, flo // 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 _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
