r365798 - [CodeGen] NVPTX: Switch from atomic.load.add.f32 to atomicrmw fadd

[Benjamin Kramer via cfe-commits]

Author: d0k
Date: Thu Jul 11 10:44:11 2019
New Revision: 365798

URL: http://llvm.org/viewvc/llvm-project?rev=365798&view=rev
Log:
[CodeGen] NVPTX: Switch from atomic.load.add.f32 to atomicrmw fadd

Modified:
    cfe/trunk/lib/CodeGen/CGBuiltin.cpp
    cfe/trunk/test/CodeGen/builtins-nvptx-ptx50.cu
    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=365798&r1=365797&r2=365798&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CGBuiltin.cpp (original)
+++ cfe/trunk/lib/CodeGen/CGBuiltin.cpp Thu Jul 11 10:44:11 2019
@@ -13472,24 +13472,12 @@ CodeGenFunction::EmitNVPTXBuiltinExpr(un
     // success flag.
     return MakeAtomicCmpXchgValue(*this, E, /*ReturnBool=*/false);
 
-  case NVPTX::BI__nvvm_atom_add_gen_f: {
-    Value *Ptr = EmitScalarExpr(E->getArg(0));
-    Value *Val = EmitScalarExpr(E->getArg(1));
-    // atomicrmw only deals with integer arguments so we need to use
-    // LLVM's nvvm_atomic_load_add_f32 intrinsic for that.
-    Function *FnALAF32 =
-        CGM.getIntrinsic(Intrinsic::nvvm_atomic_load_add_f32, Ptr->getType());
-    return Builder.CreateCall(FnALAF32, {Ptr, Val});
-  }
-
+  case NVPTX::BI__nvvm_atom_add_gen_f:
   case NVPTX::BI__nvvm_atom_add_gen_d: {
     Value *Ptr = EmitScalarExpr(E->getArg(0));
     Value *Val = EmitScalarExpr(E->getArg(1));
-    // atomicrmw only deals with integer arguments, so we need to use
-    // LLVM's nvvm_atomic_load_add_f64 intrinsic.
-    Function *FnALAF64 =
-        CGM.getIntrinsic(Intrinsic::nvvm_atomic_load_add_f64, Ptr->getType());
-    return Builder.CreateCall(FnALAF64, {Ptr, Val});
+    return Builder.CreateAtomicRMW(llvm::AtomicRMWInst::FAdd, Ptr, Val,
+                                   AtomicOrdering::SequentiallyConsistent);
   }
 
   case NVPTX::BI__nvvm_atom_inc_gen_ui: {

Modified: cfe/trunk/test/CodeGen/builtins-nvptx-ptx50.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGen/builtins-nvptx-ptx50.cu?rev=365798&r1=365797&r2=365798&view=diff
==============================================================================
--- cfe/trunk/test/CodeGen/builtins-nvptx-ptx50.cu (original)
+++ cfe/trunk/test/CodeGen/builtins-nvptx-ptx50.cu Thu Jul 11 10:44:11 2019
@@ -17,7 +17,7 @@
 
 // CHECK-LABEL: test_fn
 __device__ void test_fn(double d, double* double_ptr) {
-  // CHECK: call double @llvm.nvvm.atomic.load.add.f64.p0f64
+  // CHECK: atomicrmw fadd double
   // expected-error at +1 {{'__nvvm_atom_add_gen_d' needs target feature sm_60}}
   __nvvm_atom_add_gen_d(double_ptr, d);
 }

Modified: cfe/trunk/test/CodeGen/builtins-nvptx.c
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGen/builtins-nvptx.c?rev=365798&r1=365797&r2=365798&view=diff
==============================================================================
--- cfe/trunk/test/CodeGen/builtins-nvptx.c (original)
+++ cfe/trunk/test/CodeGen/builtins-nvptx.c Thu Jul 11 10:44:11 2019
@@ -279,7 +279,7 @@ __device__ void nvvm_atom(float *fp, flo
   // 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
+  // CHECK: atomicrmw fadd float
   __nvvm_atom_add_gen_f(fp, f);
 
   // CHECK: call i32 @llvm.nvvm.atomic.load.inc.32.p0i32