[PATCH] D18322: [CUDA] Implement atomicInc and atomicDec builtins
Justin Lebar via cfe-commits
cfe-commits at lists.llvm.org
Mon Mar 21 17:14:40 PDT 2016
This revision was automatically updated to reflect the committed changes.
Closed by commit rL264009: [CUDA] Implement atomicInc and atomicDec builtins (authored by jlebar).
Changed prior to commit:
http://reviews.llvm.org/D18322?vs=51191&id=51248#toc
Repository:
rL LLVM
http://reviews.llvm.org/D18322
Files:
cfe/trunk/lib/CodeGen/CGBuiltin.cpp
cfe/trunk/test/CodeGen/builtins-nvptx.c
Index: cfe/trunk/test/CodeGen/builtins-nvptx.c
===================================================================
--- cfe/trunk/test/CodeGen/builtins-nvptx.c
+++ cfe/trunk/test/CodeGen/builtins-nvptx.c
@@ -189,7 +189,7 @@
// Check for atomic intrinsics
// CHECK-LABEL: nvvm_atom
-__device__ void nvvm_atom(float *fp, float f, int *ip, int i, long *lp, long l,
+__device__ void nvvm_atom(float *fp, float f, int *ip, int i, unsigned int *uip, unsigned ui, long *lp, long l,
long long *llp, long long ll) {
// CHECK: atomicrmw add
__nvvm_atom_add_gen_i(ip, i);
@@ -272,5 +272,11 @@
// CHECK: call float @llvm.nvvm.atomic.load.add.f32.p0f32
__nvvm_atom_add_gen_f(fp, f);
+ // CHECK: call i32 @llvm.nvvm.atomic.load.inc.32.p0i32
+ __nvvm_atom_inc_gen_ui(uip, ui);
+
+ // CHECK: call i32 @llvm.nvvm.atomic.load.dec.32.p0i32
+ __nvvm_atom_dec_gen_ui(uip, ui);
+
// CHECK: ret
}
Index: cfe/trunk/lib/CodeGen/CGBuiltin.cpp
===================================================================
--- cfe/trunk/lib/CodeGen/CGBuiltin.cpp
+++ cfe/trunk/lib/CodeGen/CGBuiltin.cpp
@@ -7439,6 +7439,22 @@
return Builder.CreateCall(FnALAF32, {Ptr, Val});
}
+ case NVPTX::BI__nvvm_atom_inc_gen_ui: {
+ Value *Ptr = EmitScalarExpr(E->getArg(0));
+ Value *Val = EmitScalarExpr(E->getArg(1));
+ Value *FnALI32 =
+ CGM.getIntrinsic(Intrinsic::nvvm_atomic_load_inc_32, Ptr->getType());
+ return Builder.CreateCall(FnALI32, {Ptr, Val});
+ }
+
+ case NVPTX::BI__nvvm_atom_dec_gen_ui: {
+ Value *Ptr = EmitScalarExpr(E->getArg(0));
+ Value *Val = EmitScalarExpr(E->getArg(1));
+ Value *FnALD32 =
+ CGM.getIntrinsic(Intrinsic::nvvm_atomic_load_dec_32, Ptr->getType());
+ return Builder.CreateCall(FnALD32, {Ptr, Val});
+ }
+
default:
return nullptr;
}
-------------- next part --------------
A non-text attachment was scrubbed...
Name: D18322.51248.patch
Type: text/x-patch
Size: 1839 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20160322/b37eb289/attachment.bin>
More information about the cfe-commits
mailing list