OpenCL: add support for __kernel, kernel keywords and EXTENSION,

FP_CONTRACT pragmas.  Patch originally by ARM.

git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@125475 91177308-0d34-0410-b5e6-96231b3b80d8

diff --git a/include/clang/Basic/Attr.td b/include/clang/Basic/Attr.td
index 9b0982a29b9f0afb52ffd90eb99464e20b7da55a..3e62d411d51a348797b97c7ab8dddad5aaca8a8f 100644 (file)
--- a/include/clang/Basic/Attr.td
+++ b/include/clang/Basic/Attr.td
@@ -202,6 +202,10 @@ def CUDAShared : InheritableAttr {
   let Spellings = ["shared"];
 }
 
+def OpenCLKernel : Attr {
+  let Spellings = ["opencl_kernel_function"];
+}
+
 def Deprecated : InheritableAttr {
   let Spellings = ["deprecated"];
   let Args = [StringArgument<"Message">];

diff --git a/include/clang/Basic/DiagnosticParseKinds.td b/include/clang/Basic/DiagnosticParseKinds.td
index 1eedf639cecbbbb72b9075fbb9f2f1185091247b..8d4aa2868f40bc669ade79365f6fe8f868f65871 100644 (file)
--- a/include/clang/Basic/DiagnosticParseKinds.td
+++ b/include/clang/Basic/DiagnosticParseKinds.td
@@ -447,5 +447,13 @@ def warn_pragma_unused_expected_punc : Warning<
 def err_not_opencl_storage_class_specifier : Error<
   "OpenCL does not support the '%0' storage class specifier">;
 
+// OpenCL EXTENSION pragma (OpenCL 1.1 [9.1])
+def warn_pragma_expected_colon : Warning<
+  "missing ':' after %0 - ignoring">;
+def warn_pragma_expected_enable_disable : Warning<
+  "expected 'enable' or 'disable' - ignoring">;
+def warn_pragma_unknown_extension : Warning<
+  "unknown OpenCL extension %0 - ignoring">;
+
 } // end of Parse Issue category.
 } // end of Parser diagnostics

diff --git a/include/clang/Basic/LangOptions.h b/include/clang/Basic/LangOptions.h
index 6267e65fbe28ea9dc424eb952e6354610df2df33..b5877c03ccb1e15a0426605779fd4fcc07d3f0da 100644 (file)
--- a/include/clang/Basic/LangOptions.h
+++ b/include/clang/Basic/LangOptions.h
@@ -249,6 +249,18 @@ public:
     fp_contract(LangOpts.DefaultFPContract) {}
 };
 
+/// OpenCL volatile options
+class OpenCLOptions {
+public:
+#define OPENCLEXT(nm)  unsigned nm : 1;
+#include "clang/Basic/OpenCLExtensions.def"
+
+  OpenCLOptions() {
+#define OPENCLEXT(nm)   nm = 0;
+#include "clang/Basic/OpenCLExtensions.def"
+  }
+};
+
 }  // end namespace clang
 
 #endif

diff --git a/include/clang/Basic/TokenKinds.def b/include/clang/Basic/TokenKinds.def
index cf917a95cea2ff1612c5207b3d6601e84f849274..b84b04da3de2c1e3d2a4cbebfbd8895117c84989 100644 (file)
--- a/include/clang/Basic/TokenKinds.def
+++ b/include/clang/Basic/TokenKinds.def
@@ -192,6 +192,8 @@ PUNCTUATOR(greatergreatergreater, ">>>")
 //   KEYCXX0X - This is a C++ keyword introduced to C++ in C++0x
 //   KEYGNU   - This is a keyword if GNU extensions are enabled
 //   KEYMS    - This is a keyword if Microsoft extensions are enabled
+//   KEYOPENCL  - This is a keyword in OpenCL
+//   KEYALTIVEC - This is a keyword in AltiVec
 //   KEYBORLAND - This is a keyword if Borland extensions are enabled
 //
 KEYWORD(auto                        , KEYALL)
@@ -343,6 +345,10 @@ KEYWORD(__fastcall                  , KEYALL)
 KEYWORD(__thiscall                  , KEYALL)
 KEYWORD(__forceinline               , KEYALL)
 
+// OpenCL-specific keywords (see OpenCL 1.1 [6.1.9])
+KEYWORD(__kernel                    , KEYOPENCL)
+ALIAS("kernel", __kernel            , KEYOPENCL)
+
 // Borland Extensions.
 KEYWORD(__pascal                    , KEYALL)
 

diff --git a/include/clang/Parse/Parser.h b/include/clang/Parse/Parser.h
index fe65d06e6c38ecbee21ec8a3a95817099f66f60d..c67b451b875213b2cd7343e05e9803236f19b3f4 100644 (file)
--- a/include/clang/Parse/Parser.h
+++ b/include/clang/Parse/Parser.h
@@ -123,6 +123,7 @@ class Parser : public CodeCompletionHandler {
   llvm::OwningPtr<PragmaHandler> UnusedHandler;
   llvm::OwningPtr<PragmaHandler> WeakHandler;
   llvm::OwningPtr<PragmaHandler> FPContractHandler;
+  llvm::OwningPtr<PragmaHandler> OpenCLExtensionHandler;
 
   /// Whether the '>' token acts as an operator or not. This will be
   /// true except when we are parsing an expression within a C++
@@ -1526,6 +1527,7 @@ private:
   void ParseMicrosoftDeclSpec(ParsedAttributes &attrs);
   void ParseMicrosoftTypeAttributes(ParsedAttributes &attrs);
   void ParseBorlandTypeAttributes(ParsedAttributes &attrs);
+  void ParseOpenCLAttributes(ParsedAttributes &attrs);
 
   void ParseTypeofSpecifier(DeclSpec &DS);
   void ParseDecltypeSpecifier(DeclSpec &DS);

diff --git a/include/clang/Sema/AttributeList.h b/include/clang/Sema/AttributeList.h
index 91389a4d9847c3c952689f794e3fe4ea80444982..c675e75f47a5a82127ba6153063609768562e2be 100644 (file)
--- a/include/clang/Sema/AttributeList.h
+++ b/include/clang/Sema/AttributeList.h
@@ -134,6 +134,7 @@ public:
     AT_ns_consumed,             // Clang-specific.
     AT_ns_consumes_self,        // Clang-specific.
     AT_objc_gc,
+    AT_opencl_kernel_function,  // OpenCL-specific.
     AT_overloadable,       // Clang-specific.
     AT_ownership_holds,    // Clang-specific.
     AT_ownership_returns,  // Clang-specific.

diff --git a/include/clang/Sema/Sema.h b/include/clang/Sema/Sema.h
index f35fef89529264f61a618a142c79128ea941523e..f259cb9ca943115ac32038bdb0964d1693f3bb3f 100644 (file)
--- a/include/clang/Sema/Sema.h
+++ b/include/clang/Sema/Sema.h
@@ -210,6 +210,7 @@ public:
   typedef TemplateParameterList TemplateParamsTy;
   typedef NestedNameSpecifier CXXScopeTy;
 
+  OpenCLOptions OpenCLFeatures;
   FPOptions FPFeatures;
 
   const LangOptions &LangOpts;
@@ -546,6 +547,7 @@ public:
   void Initialize();
   
   const LangOptions &getLangOptions() const { return LangOpts; }
+  OpenCLOptions &getOpenCLOptions() { return OpenCLFeatures; }
   FPOptions     &getFPOptions() { return FPFeatures; }
 
   Diagnostic &getDiagnostics() const { return Diags; }
@@ -4412,7 +4414,7 @@ public:
                             SourceLocation AliasNameLoc);
 
   /// ActOnPragmaFPContract - Called on well formed
-  /// #pragma STDC FP_CONTRACT
+  /// #pragma {STDC,OPENCL} FP_CONTRACT
   void ActOnPragmaFPContract(tok::OnOffSwitch OOS);
 
   /// AddAlignmentAttributesForRecord - Adds any needed alignment attributes to

diff --git a/lib/Basic/IdentifierTable.cpp b/lib/Basic/IdentifierTable.cpp
index 47f12923a21133d35442bb547f7b0178364b9bb5..48a5f49914b2a536b1bdd6eab550250c07c9e198 100644 (file)
--- a/lib/Basic/IdentifierTable.cpp
+++ b/lib/Basic/IdentifierTable.cpp
@@ -90,7 +90,8 @@ namespace {
     BOOLSUPPORT = 64,
     KEYALTIVEC = 128,
     KEYNOCXX = 256,
-    KEYBORLAND = 512
+    KEYBORLAND = 512,
+    KEYOPENCL = 1024
   };
 }
 
@@ -115,6 +116,7 @@ static void AddKeyword(llvm::StringRef Keyword,
   else if (LangOpts.Borland && (Flags & KEYBORLAND)) AddResult = 1;
   else if (LangOpts.Bool && (Flags & BOOLSUPPORT)) AddResult = 2;
   else if (LangOpts.AltiVec && (Flags & KEYALTIVEC)) AddResult = 2;
+  else if (LangOpts.OpenCL && (Flags & KEYOPENCL)) AddResult = 2;
   else if (!LangOpts.CPlusPlus && (Flags & KEYNOCXX)) AddResult = 2;
 
   // Don't add this keyword if disabled in this language.

diff --git a/lib/CodeGen/CodeGenFunction.cpp b/lib/CodeGen/CodeGenFunction.cpp
index 38ca0214da3041973cd912323619e887c92177b6..39aff78c73c25d6ed923d27acae43ffe17e1ba2a 100644 (file)
--- a/lib/CodeGen/CodeGenFunction.cpp
+++ b/lib/CodeGen/CodeGenFunction.cpp
@@ -239,6 +239,19 @@ void CodeGenFunction::StartFunction(GlobalDecl GD, QualType RetTy,
         break;
       }
 
+  if (getContext().getLangOptions().OpenCL) {
+    // Add metadata for a kernel function.
+    if (const FunctionDecl *FD = dyn_cast_or_null<FunctionDecl>(D))
+      if (FD->hasAttr<OpenCLKernelAttr>()) {
+        llvm::LLVMContext &Context = getLLVMContext();
+        llvm::NamedMDNode *OpenCLMetadata = 
+          CGM.getModule().getOrInsertNamedMetadata("opencl.kernels");
+          
+        llvm::Value *Op = Fn;
+        OpenCLMetadata->addOperand(llvm::MDNode::get(Context, &Op, 1));
+      }
+  }
+
   llvm::BasicBlock *EntryBB = createBasicBlock("entry", CurFn);
 
   // Create a marker to make it easy to insert allocas into the entryblock

diff --git a/lib/Frontend/CompilerInvocation.cpp b/lib/Frontend/CompilerInvocation.cpp
index 21038565e1fd4262bb9500043402885174dd373c..229ea504abb3ff9d2dc1268751a5d5e3da60399e 100644 (file)
--- a/lib/Frontend/CompilerInvocation.cpp
+++ b/lib/Frontend/CompilerInvocation.cpp
@@ -1333,6 +1333,7 @@ void CompilerInvocation::setLangDefaults(LangOptions &Opts, InputKind IK,
     Opts.AltiVec = 1;
     Opts.CXXOperatorNames = 1;
     Opts.LaxVectorConversions = 1;
+    Opts.DefaultFPContract = 1;
   }
 
   if (LangStd == LangStandard::lang_cuda)

diff --git a/lib/Parse/ParseDecl.cpp b/lib/Parse/ParseDecl.cpp
index d97b4e30a0ae74211b53fb5abf449ea140cbb224..b815031939833e082e196b6dd38c7c281593e651 100644 (file)
--- a/lib/Parse/ParseDecl.cpp
+++ b/lib/Parse/ParseDecl.cpp
@@ -301,6 +301,16 @@ void Parser::ParseBorlandTypeAttributes(ParsedAttributes &attrs) {
   }
 }
 
+void Parser::ParseOpenCLAttributes(ParsedAttributes &attrs) {
+  // Treat these like attributes
+  while (Tok.is(tok::kw___kernel)) {
+    SourceLocation AttrNameLoc = ConsumeToken();
+    attrs.add(AttrFactory.Create(PP.getIdentifierInfo("opencl_kernel_function"),
+                                 AttrNameLoc, 0, AttrNameLoc, 0,
+                                 SourceLocation(), 0, 0, false));
+  }
+}
+
 void Parser::DiagnoseProhibitedAttributes(ParsedAttributesWithRange &attrs) {
   Diag(attrs.Range.getBegin(), diag::err_attributes_not_allowed)
     << attrs.Range;
@@ -864,6 +874,7 @@ Parser::getDeclSpecContextFromDeclaratorContext(unsigned Context) {
 /// [C99]   'inline'
 /// [C++]   'virtual'
 /// [C++]   'explicit'
+/// [OpenCL] '__kernel'
 ///       'friend': [C++ dcl.friend]
 ///       'constexpr': [C++0x dcl.constexpr]
 
@@ -1201,6 +1212,11 @@ void Parser::ParseDeclarationSpecifiers(DeclSpec &DS,
       ParseBorlandTypeAttributes(DS.getAttributes());
       continue;
 
+    // OpenCL single token adornments.
+    case tok::kw___kernel:
+      ParseOpenCLAttributes(DS.getAttributes());
+      continue;
+
     // storage-class-specifier
     case tok::kw_typedef:
       isInvalid = DS.SetStorageClassSpec(DeclSpec::SCS_typedef, Loc, PrevSpec,

diff --git a/lib/Parse/ParsePragma.cpp b/lib/Parse/ParsePragma.cpp
index 41f32fb945662d6e81fc24d4817acb710618a45c..dfd0da079df5d3ec427e4eb223e7b465668fa50b 100644 (file)
--- a/lib/Parse/ParsePragma.cpp
+++ b/lib/Parse/ParsePragma.cpp
@@ -382,3 +382,53 @@ PragmaFPContractHandler::HandlePragma(Preprocessor &PP,
 
   Actions.ActOnPragmaFPContract(OOS);
 }
+
+void 
+PragmaOpenCLExtensionHandler::HandlePragma(Preprocessor &PP, 
+                                           PragmaIntroducerKind Introducer,
+                                           Token &Tok) {
+  PP.Lex(Tok);
+  if (Tok.isNot(tok::identifier)) {
+    PP.Diag(Tok.getLocation(), diag::warn_pragma_expected_identifier) <<
+      "OPENCL";
+    return;
+  }
+  IdentifierInfo *ename = Tok.getIdentifierInfo();
+  SourceLocation NameLoc = Tok.getLocation();
+
+  PP.Lex(Tok);
+  if (Tok.isNot(tok::colon)) {
+    PP.Diag(Tok.getLocation(), diag::warn_pragma_expected_colon) << ename;
+    return;
+  }
+
+  PP.Lex(Tok);
+  if (Tok.isNot(tok::identifier)) {
+    PP.Diag(Tok.getLocation(), diag::warn_pragma_expected_enable_disable);
+    return;
+  }
+  IdentifierInfo *op = Tok.getIdentifierInfo();
+
+  unsigned state;
+  if (op->isStr("enable")) {
+    state = 1;
+  } else if (op->isStr("disable")) {
+    state = 0;
+  } else {
+    PP.Diag(Tok.getLocation(), diag::warn_pragma_expected_enable_disable);
+    return;
+  }
+
+  OpenCLOptions &f = Actions.getOpenCLOptions();
+  if (ename->isStr("all")) {
+#define OPENCLEXT(nm)   f.nm = state;
+#include "clang/Basic/OpenCLExtensions.def"
+  }
+#define OPENCLEXT(nm) else if (ename->isStr(#nm)) { f.nm = state; }
+#include "clang/Basic/OpenCLExtensions.def"
+  else {
+    PP.Diag(NameLoc, diag::warn_pragma_unknown_extension) << ename;
+    return;
+  }
+}
+

diff --git a/lib/Parse/ParsePragma.h b/lib/Parse/ParsePragma.h
index 80894b28d85dec1379267705561377dee05f76fe..bee6af3f4cfddef17188725e238c0c13b25ed965 100644 (file)
--- a/lib/Parse/ParsePragma.h
+++ b/lib/Parse/ParsePragma.h
@@ -80,6 +80,17 @@ public:
                             Token &FirstToken);
 };
 
+class PragmaOpenCLExtensionHandler : public PragmaHandler {
+  Sema &Actions;
+  Parser &parser;
+public:
+  PragmaOpenCLExtensionHandler(Sema &S, Parser& p) : 
+    PragmaHandler("EXTENSION"), Actions(S), parser(p) {}
+  virtual void HandlePragma(Preprocessor &PP, PragmaIntroducerKind Introducer,
+                            Token &FirstToken);
+};
+  
+
 class PragmaFPContractHandler : public PragmaHandler {
   Sema &Actions;
   Parser &parser;

diff --git a/lib/Parse/Parser.cpp b/lib/Parse/Parser.cpp
index 8273d5e2d15a8354829d7e5312956b780fef8c35..a50763a0e38e146dbdccbe8b2bc4bb62fd505613 100644 (file)
--- a/lib/Parse/Parser.cpp
+++ b/lib/Parse/Parser.cpp
@@ -53,6 +53,14 @@ Parser::Parser(Preprocessor &pp, Sema &actions)
 
   FPContractHandler.reset(new PragmaFPContractHandler(actions, *this));
   PP.AddPragmaHandler("STDC", FPContractHandler.get());
+
+  if (getLang().OpenCL) {
+    OpenCLExtensionHandler.reset(
+                  new PragmaOpenCLExtensionHandler(actions, *this));
+    PP.AddPragmaHandler("OPENCL", OpenCLExtensionHandler.get());
+
+    PP.AddPragmaHandler("OPENCL", FPContractHandler.get());
+  }
       
   PP.setCodeCompletionHandler(*this);
 }
@@ -363,6 +371,13 @@ Parser::~Parser() {
   UnusedHandler.reset();
   PP.RemovePragmaHandler(WeakHandler.get());
   WeakHandler.reset();
+
+  if (getLang().OpenCL) {
+    PP.RemovePragmaHandler("OPENCL", OpenCLExtensionHandler.get());
+    OpenCLExtensionHandler.reset();
+    PP.RemovePragmaHandler("OPENCL", FPContractHandler.get());
+  }
+
   PP.RemovePragmaHandler("STDC", FPContractHandler.get());
   FPContractHandler.reset();
   PP.clearCodeCompletionHandler();

diff --git a/lib/Sema/AttributeList.cpp b/lib/Sema/AttributeList.cpp
index 77d962542bfc47b34b81468b2b0a5fd87aafe8ab..73e956cc91a4db77db3d85a933d493e2e03b484c 100644 (file)
--- a/lib/Sema/AttributeList.cpp
+++ b/lib/Sema/AttributeList.cpp
@@ -134,6 +134,7 @@ AttributeList::Kind AttributeList::getKind(const IdentifierInfo *Name) {
     .Case("launch_bounds", AT_launch_bounds)
     .Case("common", AT_common)
     .Case("nocommon", AT_nocommon)
+    .Case("opencl_kernel_function", AT_opencl_kernel_function)
     .Case("uuid", AT_uuid)
     .Default(UnknownAttribute);
 }

diff --git a/lib/Sema/SemaDeclAttr.cpp b/lib/Sema/SemaDeclAttr.cpp
index 3d64ade17304dd40c10f78a8df092643d4a7618b..5d5093f5fe42804f17d4d29a2d7b021b5429b51e 100644 (file)
--- a/lib/Sema/SemaDeclAttr.cpp
+++ b/lib/Sema/SemaDeclAttr.cpp
@@ -2333,6 +2333,11 @@ static void HandleCallConvAttr(Decl *d, const AttributeList &attr, Sema &S) {
   }
 }
 
+static void HandleOpenCLKernelAttr(Decl *d, const AttributeList &Attr, Sema &S){
+  assert(Attr.isInvalid() == false);
+  d->addAttr(::new (S.Context) OpenCLKernelAttr(Attr.getLoc(), S.Context));
+}
+
 bool Sema::CheckCallingConvAttr(const AttributeList &attr, CallingConv &CC) {
   if (attr.isInvalid())
     return true;
@@ -2774,6 +2779,9 @@ static void ProcessInheritableDeclAttr(Scope *scope, Decl *D,
   case AttributeList::AT_pascal:
     HandleCallConvAttr(D, Attr, S);
     break;
+  case AttributeList::AT_opencl_kernel_function:
+    HandleOpenCLKernelAttr(D, Attr, S);
+    break;
   case AttributeList::AT_uuid:
     HandleUuidAttr(D, Attr, S);
     break;

diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt
index 4c15e8f5f8624c61518e69f7b22b8f31777c7531..e16416a34b196068df657cc8fc01cde43afd8721 100644 (file)
--- a/test/CMakeLists.txt
+++ b/test/CMakeLists.txt
@@ -4,6 +4,7 @@ set(CLANG_TEST_DIRECTORIES
   "CodeGen"
   "CodeGenCXX"
   "CodeGenObjC"
+  "CodeGenOpenCL"
   "Coverage"
   "CXX"
   "Driver"

diff --git a/test/CodeGenOpenCL/kernel-metadata.cl b/test/CodeGenOpenCL/kernel-metadata.cl
new file mode 100644 (file)
index 0000000..3e10a11
--- /dev/null
+++ b/test/CodeGenOpenCL/kernel-metadata.cl
@@ -0,0 +1,10 @@
+// RUN: %clang_cc1 %s -emit-llvm -o - | FileCheck %s
+
+void normal_function() {
+}
+
+__kernel void kernel_function() {
+}
+
+// CHECK: !opencl.kernels = !{!0}
+// CHECK: !0 = metadata !{void ()* @kernel_function}

diff --git a/test/Parser/opencl-kernel.cl b/test/Parser/opencl-kernel.cl
new file mode 100644 (file)
index 0000000..3abb62b
--- /dev/null
+++ b/test/Parser/opencl-kernel.cl
@@ -0,0 +1,9 @@
+// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only
+
+__kernel void test()
+{
+}
+
+kernel void test1()
+{
+}

diff --git a/test/Parser/opencl-pragma.cl b/test/Parser/opencl-pragma.cl
new file mode 100644 (file)
index 0000000..5b6c55a
--- /dev/null
+++ b/test/Parser/opencl-pragma.cl
@@ -0,0 +1,12 @@
+// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only
+
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+
+#pragma OPENCL EXTENSION cl_no_such_extension : disable /* expected-warning {{unknown OpenCL extension 'cl_no_such_extension' - ignoring}} */
+
+#pragma OPENCL EXTENSION cl_khr_fp64 : on /* expected-warning {{expected 'enable' or 'disable' - ignoring}} */
+
+#pragma OPENCL FP_CONTRACT ON
+#pragma OPENCL FP_CONTRACT OFF
+#pragma OPENCL FP_CONTRACT DEFAULT
+#pragma OPENCL FP_CONTRACT FOO // expected-warning {{expected 'ON' or 'OFF' or 'DEFAULT' in pragma}}