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

[Tom Stellard]

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

-------------- 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
-------------- next part --------------
diff --git include/llvm/Intrinsics.td include/llvm/Intrinsics.td
index 069f907..c1af0cf 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_i32_ty], [llvm_i32_ty]>;
+
+
 //===----------------------------------------------------------------------===//
 // Target-specific intrinsics
 //===----------------------------------------------------------------------===//