From 38c658f9ae1407a9d44e864573b31250eb8ca8af Mon Sep 17 00:00:00 2001 From: Benjamin Kramer Date: Thu, 11 Jul 2019 17:44:11 +0000 Subject: [PATCH] [CodeGen] NVPTX: Switch from atomic.load.add.f32 to atomicrmw fadd git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@365798 91177308-0d34-0410-b5e6-96231b3b80d8 --- lib/CodeGen/CGBuiltin.cpp | 18 +++--------------- test/CodeGen/builtins-nvptx-ptx50.cu | 2 +- test/CodeGen/builtins-nvptx.c | 2 +- 3 files changed, 5 insertions(+), 17 deletions(-) diff --git a/lib/CodeGen/CGBuiltin.cpp b/lib/CodeGen/CGBuiltin.cpp index 52e2d5bfb9..086785fdba 100644 --- a/lib/CodeGen/CGBuiltin.cpp +++ b/lib/CodeGen/CGBuiltin.cpp @@ -13472,24 +13472,12 @@ CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID, const CallExpr *E) { // 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: { diff --git a/test/CodeGen/builtins-nvptx-ptx50.cu b/test/CodeGen/builtins-nvptx-ptx50.cu index 72e1aecb48..4436ff523c 100644 --- a/test/CodeGen/builtins-nvptx-ptx50.cu +++ b/test/CodeGen/builtins-nvptx-ptx50.cu @@ -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@+1 {{'__nvvm_atom_add_gen_d' needs target feature sm_60}} __nvvm_atom_add_gen_d(double_ptr, d); } diff --git a/test/CodeGen/builtins-nvptx.c b/test/CodeGen/builtins-nvptx.c index 16f41bac34..31c3ecdb14 100644 --- a/test/CodeGen/builtins-nvptx.c +++ b/test/CodeGen/builtins-nvptx.c @@ -279,7 +279,7 @@ __device__ void nvvm_atom(float *fp, float f, double *dfp, double df, int *ip, // 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 -- 2.40.0