aboutsummaryrefslogtreecommitdiff
diff options
context:
space:
mode:
Diffstat
-rw-r--r--include/clang/Basic/Attr.td4
-rw-r--r--include/clang/Basic/DiagnosticParseKinds.td8
-rw-r--r--include/clang/Basic/LangOptions.h12
-rw-r--r--include/clang/Basic/TokenKinds.def6
-rw-r--r--include/clang/Parse/Parser.h2
-rw-r--r--include/clang/Sema/AttributeList.h1
-rw-r--r--include/clang/Sema/Sema.h4
-rw-r--r--lib/Basic/IdentifierTable.cpp4
-rw-r--r--lib/CodeGen/CodeGenFunction.cpp13
-rw-r--r--lib/Frontend/CompilerInvocation.cpp1
-rw-r--r--lib/Parse/ParseDecl.cpp16
-rw-r--r--lib/Parse/ParsePragma.cpp50
-rw-r--r--lib/Parse/ParsePragma.h11
-rw-r--r--lib/Parse/Parser.cpp15
-rw-r--r--lib/Sema/AttributeList.cpp1
-rw-r--r--lib/Sema/SemaDeclAttr.cpp8
-rw-r--r--test/CMakeLists.txt1
-rw-r--r--test/CodeGenOpenCL/kernel-metadata.cl10
-rw-r--r--test/Parser/opencl-kernel.cl9
-rw-r--r--test/Parser/opencl-pragma.cl12
20 files changed, 186 insertions, 2 deletions
diff --git a/include/clang/Basic/Attr.td b/include/clang/Basic/Attr.td
index 9b0982a29b..3e62d411d5 100644
--- 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 1eedf639ce..8d4aa2868f 100644
--- 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 6267e65fbe..b5877c03cc 100644
--- 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 cf917a95ce..b84b04da3d 100644
--- 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 fe65d06e6c..c67b451b87 100644
--- 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 91389a4d98..c675e75f47 100644
--- 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 f35fef8952..f259cb9ca9 100644
--- 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 47f12923a2..48a5f49914 100644
--- 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 38ca0214da..39aff78c73 100644
--- 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 21038565e1..229ea504ab 100644
--- 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 d97b4e30a0..b815031939 100644
--- 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 41f32fb945..dfd0da079d 100644
--- 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 80894b28d8..bee6af3f4c 100644
--- 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 8273d5e2d1..a50763a0e3 100644
--- 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 77d962542b..73e956cc91 100644
--- 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 3d64ade173..5d5093f5fe 100644
--- 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 4c15e8f5f8..e16416a34b 100644
--- 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
index 0000000000..3e10a119d0
--- /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
index 0000000000..3abb62b616
--- /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
index 0000000000..5b6c55ab07
--- /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}}
generated by cgit v1.2.3-18-g5258 (git 2.32.0) at 2025-07-02 22:54:36 +0000