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@+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
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
