[PATCH] D12487: [CUDA] fix codegen for __nvvm_atom_min/max_gen_u*

[Jingyue Wu via llvm-commits]

jingyue created this revision.
jingyue added reviewers: tra, eliben.
jingyue added a subscriber: llvm-commits.

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

http://reviews.llvm.org/D12487

Files:
  lib/CodeGen/CGBuiltin.cpp
  test/CodeGenCUDA/atomic.cu

Index: test/CodeGenCUDA/atomic.cu
===================================================================
--- /dev/null
+++ test/CodeGenCUDA/atomic.cu
@@ -0,0 +1,45 @@
+// RUN: %clang_cc1 -emit-llvm %s -o - -fcuda-is-device -triple nvptx-unknown-unknown | FileCheck %s
+
+#include "Inputs/cuda.h"
+
+__device__ void TestMin(int i, int *min_i,
+                        long l, long *min_l,
+                        long long ll, long long *min_ll,
+                        unsigned int ui, unsigned int *min_ui,
+                        unsigned long ul, unsigned long *min_ul,
+                        unsigned long long ull, unsigned long long *min_ull) {
+  // CHECK-LABEL: define void @_Z7TestMiniPilPlxPxjPjmPmyPy(
+  __nvvm_atom_min_gen_i(min_i, i);
+  // CHECK: atomicrmw min i32*
+  __nvvm_atom_min_gen_l(min_l, l);
+  // CHECK: atomicrmw min i64*
+  __nvvm_atom_min_gen_ll(min_ll, ll);
+  // CHECK: atomicrmw min i64*
+  __nvvm_atom_min_gen_ui(min_ui, ui);
+  // CHECK: atomicrmw umin i32*
+  __nvvm_atom_min_gen_ul(min_ul, ul);
+  // CHECK: atomicrmw umin i64*
+  __nvvm_atom_min_gen_ull(min_ull, ull);
+  // CHECK: atomicrmw umin i64*
+}
+
+__device__ void TestMax(int i, int *max_i,
+                        long l, long *max_l,
+                        long long ll, long long *max_ll,
+                        unsigned int ui, unsigned int *max_ui,
+                        unsigned long ul, unsigned long *max_ul,
+                        unsigned long long ull, unsigned long long *max_ull) {
+  // CHECK-LABEL: define void @_Z7TestMaxiPilPlxPxjPjmPmyPy(
+  __nvvm_atom_max_gen_i(max_i, i);
+  // CHECK: atomicrmw max i32*
+  __nvvm_atom_max_gen_l(max_l, l);
+  // CHECK: atomicrmw max i64*
+  __nvvm_atom_max_gen_ll(max_ll, ll);
+  // CHECK: atomicrmw max i64*
+  __nvvm_atom_max_gen_ui(max_ui, ui);
+  // CHECK: atomicrmw umax i32*
+  __nvvm_atom_max_gen_ul(max_ul, ul);
+  // CHECK: atomicrmw umax i64*
+  __nvvm_atom_max_gen_ull(max_ull, ull);
+  // CHECK: atomicrmw umax i64*
+}
Index: lib/CodeGen/CGBuiltin.cpp
===================================================================
--- lib/CodeGen/CGBuiltin.cpp
+++ lib/CodeGen/CGBuiltin.cpp
@@ -7024,18 +7024,22 @@
   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:


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D12487.33551.patch
Type: text/x-patch
Size: 3186 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/llvm-commits/attachments/20150830/bf471993/attachment.bin>