[CUDA] fix codegen for __nvvm_atom_min/max_gen_u*

[jingyue]

Summary

Clang should emit "atomicrmw umin/umax" instead of "atomicrmw min/max".

[jingyue]

use existing tests

[tra]

LGTM.

lib/CodeGen/CGBuiltin.cpp

Value *CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID,
   case NVPTX::BI__nvvm_atom_xchg_gen_i:
   case NVPTX::BI__nvvm_atom_xchg_gen_l:
   case NVPTX::BI__nvvm_atom_xchg_gen_ll:
     return MakeBinaryAtomicValue(*this, llvm::AtomicRMWInst::Xchg, E);
 
   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:
   case NVPTX::BI__nvvm_atom_cas_gen_ll:
     return MakeAtomicCmpXchgValue(*this, E, true);
 
   case NVPTX::BI__nvvm_atom_add_gen_f: {
     Value *Ptr = EmitScalarExpr(E->getArg(0));

test/CodeGen/builtins-nvptx.c

__device__ void nvvm_atom(float *fp, float f, int *ip, int i, long *lp, long l,
 
   // CHECK: atomicrmw xchg
   __nvvm_atom_xchg_gen_i(ip, i);
   // CHECK: atomicrmw xchg
   __nvvm_atom_xchg_gen_l(&dl, 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
   __nvvm_atom_cas_gen_i(ip, 0, i);
   // CHECK: cmpxchg
   __nvvm_atom_cas_gen_l(&dl, 0, l);
   // CHECK: cmpxchg
   __nvvm_atom_cas_gen_ll(&sll, 0, ll);
 
   // CHECK: call float @llvm.nvvm.atomic.load.add.f32.p0f32
   __nvvm_atom_add_gen_f(fp, f);
 
   // CHECK: ret
 }