[cfe-commits] [llvm-commits] Patch: Add get_global_id builtin/intrinsic

[Tom Stellard]

On Fri, Feb 03, 2012 at 09:26:14PM +0000, Rotem, Nadav wrote:
> Tom, 
> 
> I can see a number of problems with this patch.  First, get_global_id does not return llvm_i32_ty on 64bit systems. It returns size_t (which is i64 on x86_64).  Second, I am not sure that this is the correct approach for implementing this. What's wrong with having get_global_id as a standard library call ? Why do we need a special intrinsic for it ?

Hi,

I've updated these two patches by changing the return type of the
get_global_id intrinsic from llvm_i32_ty to llvm_anyint_ty and the return
type of the builtin from int to size_t.  I think I still may be missing
something in Clang to make get_global_id return the correct integer
type, because the generated LLVM IR still has i32 as the return type for
get_global_id on my x86_64 system.  Hopefully, someone can spot my
mistake.

Please review.

Thanks,
Tom


> 
> Thanks,
> Nadav 
> 
> -----Original Message-----
> From: llvm-commits-bounces at cs.uiuc.edu [mailto:llvm-commits-bounces at cs.uiuc.edu] On Behalf Of Tom Stellard
> Sent: Friday, February 03, 2012 23:19
> To: llvm-commits at cs.uiuc.edu; cfe-commits at cs.uiuc.edu
> Subject: [llvm-commits] Patch: Add get_global_id builtin/intrinsic
> 
> Hi,
> 
> I've attached two patches, one for llvm and one for clang that add support for the OpenCL C builtin function get_global_id().  I would like to eventually add support for more OpenCL builtins to clang/llvm, but this initial patch is just to make sure I'm doing it the right way.
> 
> Please review.
> 
> Thanks,
> Tom Stellard
> 
> ---------------------------------------------------------------------
> Intel Israel (74) Limited
> 
> This e-mail and any attachments may contain confidential material for
> the sole use of the intended recipient(s). Any review or distribution
> by others is strictly prohibited. If you are not the intended
> recipient, please contact the sender and delete all copies.
> 
> 
-------------- next part --------------
diff --git include/llvm/Intrinsics.td include/llvm/Intrinsics.td
index 069f907..72ea8ae 100644
--- include/llvm/Intrinsics.td
+++ include/llvm/Intrinsics.td
@@ -430,6 +430,11 @@ def int_convertus  : Intrinsic<[llvm_anyint_ty],
 def int_convertuu  : Intrinsic<[llvm_anyint_ty],
                                [llvm_anyint_ty, llvm_i32_ty, llvm_i32_ty]>;
 
+//===--------------------------- OpenCL Intrinsics ------------------------===//
+//
+def int_get_global_id : Intrinsic<[llvm_anyint_ty], [llvm_i32_ty]>;
+
+
 //===----------------------------------------------------------------------===//
 // Target-specific intrinsics
 //===----------------------------------------------------------------------===//
-------------- next part --------------
diff --git include/clang/Basic/Builtins.def include/clang/Basic/Builtins.def
index b3b3c21..7e9542a 100644
--- include/clang/Basic/Builtins.def
+++ include/clang/Basic/Builtins.def
@@ -624,6 +624,8 @@ BUILTIN(__assume, "vb", "n")
 BUILTIN(__noop, "v.", "n")
 BUILTIN(__debugbreak, "v", "n")
 
+// OpenCL builtins.
+BUILTIN(__builtin_get_global_id, "ii", "n")
 
 // C99 library functions
 // C99 stdlib.h
diff --git lib/CodeGen/CGBuiltin.cpp lib/CodeGen/CGBuiltin.cpp
index e463230..d830dd0 100644
--- lib/CodeGen/CGBuiltin.cpp
+++ lib/CodeGen/CGBuiltin.cpp
@@ -1087,6 +1087,15 @@ RValue CodeGenFunction::EmitBuiltinExpr(const FunctionDecl *FD,
     llvm::StringRef Str = cast<StringLiteral>(AnnotationStrExpr)->getString();
     return RValue::get(EmitAnnotationCall(F, AnnVal, Str, E->getExprLoc()));
   }
+
+  // OpenCL builtins
+  case Builtin::BI__builtin_get_global_id: {
+    Value *Dim = EmitScalarExpr(E->getArg(0));
+    llvm::Type *ArgType = Dim->getType();
+    Value *Call = CGM.getIntrinsic(Intrinsic::get_global_id, ArgType);
+    return RValue::get(Builder.CreateCall(Call, Dim, ""));
+  }
+
   }
 
   // If this is an alias for a lib function (e.g. __builtin_sin), emit