[cfe-commits] r125475 - in /cfe/trunk: include/clang/Basic/ include/clang/Parse/ include/clang/Sema/ lib/Basic/ lib/CodeGen/ lib/Frontend/ lib/Parse/ lib/Sema/ test/ test/CodeGenOpenCL/ test/Parser/
Peter Collingbourne
peter at pcc.me.uk
Sun Feb 13 17:42:54 PST 2011
Author: pcc
Date: Sun Feb 13 19:42:53 2011
New Revision: 125475
URL: http://llvm.org/viewvc/llvm-project?rev=125475&view=rev
Log:
OpenCL: add support for __kernel, kernel keywords and EXTENSION,
FP_CONTRACT pragmas. Patch originally by ARM.
Added:
cfe/trunk/test/CodeGenOpenCL/
cfe/trunk/test/CodeGenOpenCL/kernel-metadata.cl
cfe/trunk/test/Parser/opencl-kernel.cl
cfe/trunk/test/Parser/opencl-pragma.cl
Modified:
cfe/trunk/include/clang/Basic/Attr.td
cfe/trunk/include/clang/Basic/DiagnosticParseKinds.td
cfe/trunk/include/clang/Basic/LangOptions.h
cfe/trunk/include/clang/Basic/TokenKinds.def
cfe/trunk/include/clang/Parse/Parser.h
cfe/trunk/include/clang/Sema/AttributeList.h
cfe/trunk/include/clang/Sema/Sema.h
cfe/trunk/lib/Basic/IdentifierTable.cpp
cfe/trunk/lib/CodeGen/CodeGenFunction.cpp
cfe/trunk/lib/Frontend/CompilerInvocation.cpp
cfe/trunk/lib/Parse/ParseDecl.cpp
cfe/trunk/lib/Parse/ParsePragma.cpp
cfe/trunk/lib/Parse/ParsePragma.h
cfe/trunk/lib/Parse/Parser.cpp
cfe/trunk/lib/Sema/AttributeList.cpp
cfe/trunk/lib/Sema/SemaDeclAttr.cpp
cfe/trunk/test/CMakeLists.txt
Modified: cfe/trunk/include/clang/Basic/Attr.td
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/Attr.td?rev=125475&r1=125474&r2=125475&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/Attr.td (original)
+++ cfe/trunk/include/clang/Basic/Attr.td Sun Feb 13 19:42:53 2011
@@ -202,6 +202,10 @@
let Spellings = ["shared"];
}
+def OpenCLKernel : Attr {
+ let Spellings = ["opencl_kernel_function"];
+}
+
def Deprecated : InheritableAttr {
let Spellings = ["deprecated"];
let Args = [StringArgument<"Message">];
Modified: cfe/trunk/include/clang/Basic/DiagnosticParseKinds.td
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/DiagnosticParseKinds.td?rev=125475&r1=125474&r2=125475&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/DiagnosticParseKinds.td (original)
+++ cfe/trunk/include/clang/Basic/DiagnosticParseKinds.td Sun Feb 13 19:42:53 2011
@@ -447,5 +447,13 @@
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
Modified: cfe/trunk/include/clang/Basic/LangOptions.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/LangOptions.h?rev=125475&r1=125474&r2=125475&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/LangOptions.h (original)
+++ cfe/trunk/include/clang/Basic/LangOptions.h Sun Feb 13 19:42:53 2011
@@ -249,6 +249,18 @@
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
Modified: cfe/trunk/include/clang/Basic/TokenKinds.def
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/TokenKinds.def?rev=125475&r1=125474&r2=125475&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/TokenKinds.def (original)
+++ cfe/trunk/include/clang/Basic/TokenKinds.def Sun Feb 13 19:42:53 2011
@@ -192,6 +192,8 @@
// 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(__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)
Modified: cfe/trunk/include/clang/Parse/Parser.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Parse/Parser.h?rev=125475&r1=125474&r2=125475&view=diff
==============================================================================
--- cfe/trunk/include/clang/Parse/Parser.h (original)
+++ cfe/trunk/include/clang/Parse/Parser.h Sun Feb 13 19:42:53 2011
@@ -123,6 +123,7 @@
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 @@
void ParseMicrosoftDeclSpec(ParsedAttributes &attrs);
void ParseMicrosoftTypeAttributes(ParsedAttributes &attrs);
void ParseBorlandTypeAttributes(ParsedAttributes &attrs);
+ void ParseOpenCLAttributes(ParsedAttributes &attrs);
void ParseTypeofSpecifier(DeclSpec &DS);
void ParseDecltypeSpecifier(DeclSpec &DS);
Modified: cfe/trunk/include/clang/Sema/AttributeList.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Sema/AttributeList.h?rev=125475&r1=125474&r2=125475&view=diff
==============================================================================
--- cfe/trunk/include/clang/Sema/AttributeList.h (original)
+++ cfe/trunk/include/clang/Sema/AttributeList.h Sun Feb 13 19:42:53 2011
@@ -134,6 +134,7 @@
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.
Modified: cfe/trunk/include/clang/Sema/Sema.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Sema/Sema.h?rev=125475&r1=125474&r2=125475&view=diff
==============================================================================
--- cfe/trunk/include/clang/Sema/Sema.h (original)
+++ cfe/trunk/include/clang/Sema/Sema.h Sun Feb 13 19:42:53 2011
@@ -210,6 +210,7 @@
typedef TemplateParameterList TemplateParamsTy;
typedef NestedNameSpecifier CXXScopeTy;
+ OpenCLOptions OpenCLFeatures;
FPOptions FPFeatures;
const LangOptions &LangOpts;
@@ -546,6 +547,7 @@
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 @@
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
Modified: cfe/trunk/lib/Basic/IdentifierTable.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Basic/IdentifierTable.cpp?rev=125475&r1=125474&r2=125475&view=diff
==============================================================================
--- cfe/trunk/lib/Basic/IdentifierTable.cpp (original)
+++ cfe/trunk/lib/Basic/IdentifierTable.cpp Sun Feb 13 19:42:53 2011
@@ -90,7 +90,8 @@
BOOLSUPPORT = 64,
KEYALTIVEC = 128,
KEYNOCXX = 256,
- KEYBORLAND = 512
+ KEYBORLAND = 512,
+ KEYOPENCL = 1024
};
}
@@ -115,6 +116,7 @@
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.
Modified: cfe/trunk/lib/CodeGen/CodeGenFunction.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CodeGenFunction.cpp?rev=125475&r1=125474&r2=125475&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CodeGenFunction.cpp (original)
+++ cfe/trunk/lib/CodeGen/CodeGenFunction.cpp Sun Feb 13 19:42:53 2011
@@ -239,6 +239,19 @@
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
Modified: cfe/trunk/lib/Frontend/CompilerInvocation.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Frontend/CompilerInvocation.cpp?rev=125475&r1=125474&r2=125475&view=diff
==============================================================================
--- cfe/trunk/lib/Frontend/CompilerInvocation.cpp (original)
+++ cfe/trunk/lib/Frontend/CompilerInvocation.cpp Sun Feb 13 19:42:53 2011
@@ -1333,6 +1333,7 @@
Opts.AltiVec = 1;
Opts.CXXOperatorNames = 1;
Opts.LaxVectorConversions = 1;
+ Opts.DefaultFPContract = 1;
}
if (LangStd == LangStandard::lang_cuda)
Modified: cfe/trunk/lib/Parse/ParseDecl.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Parse/ParseDecl.cpp?rev=125475&r1=125474&r2=125475&view=diff
==============================================================================
--- cfe/trunk/lib/Parse/ParseDecl.cpp (original)
+++ cfe/trunk/lib/Parse/ParseDecl.cpp Sun Feb 13 19:42:53 2011
@@ -301,6 +301,16 @@
}
}
+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 @@
/// [C99] 'inline'
/// [C++] 'virtual'
/// [C++] 'explicit'
+/// [OpenCL] '__kernel'
/// 'friend': [C++ dcl.friend]
/// 'constexpr': [C++0x dcl.constexpr]
@@ -1201,6 +1212,11 @@
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,
Modified: cfe/trunk/lib/Parse/ParsePragma.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Parse/ParsePragma.cpp?rev=125475&r1=125474&r2=125475&view=diff
==============================================================================
--- cfe/trunk/lib/Parse/ParsePragma.cpp (original)
+++ cfe/trunk/lib/Parse/ParsePragma.cpp Sun Feb 13 19:42:53 2011
@@ -382,3 +382,53 @@
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;
+ }
+}
+
Modified: cfe/trunk/lib/Parse/ParsePragma.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Parse/ParsePragma.h?rev=125475&r1=125474&r2=125475&view=diff
==============================================================================
--- cfe/trunk/lib/Parse/ParsePragma.h (original)
+++ cfe/trunk/lib/Parse/ParsePragma.h Sun Feb 13 19:42:53 2011
@@ -80,6 +80,17 @@
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;
Modified: cfe/trunk/lib/Parse/Parser.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Parse/Parser.cpp?rev=125475&r1=125474&r2=125475&view=diff
==============================================================================
--- cfe/trunk/lib/Parse/Parser.cpp (original)
+++ cfe/trunk/lib/Parse/Parser.cpp Sun Feb 13 19:42:53 2011
@@ -53,6 +53,14 @@
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 @@
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();
Modified: cfe/trunk/lib/Sema/AttributeList.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/AttributeList.cpp?rev=125475&r1=125474&r2=125475&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/AttributeList.cpp (original)
+++ cfe/trunk/lib/Sema/AttributeList.cpp Sun Feb 13 19:42:53 2011
@@ -134,6 +134,7 @@
.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);
}
Modified: cfe/trunk/lib/Sema/SemaDeclAttr.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaDeclAttr.cpp?rev=125475&r1=125474&r2=125475&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaDeclAttr.cpp (original)
+++ cfe/trunk/lib/Sema/SemaDeclAttr.cpp Sun Feb 13 19:42:53 2011
@@ -2333,6 +2333,11 @@
}
}
+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 @@
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;
Modified: cfe/trunk/test/CMakeLists.txt
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CMakeLists.txt?rev=125475&r1=125474&r2=125475&view=diff
==============================================================================
--- cfe/trunk/test/CMakeLists.txt (original)
+++ cfe/trunk/test/CMakeLists.txt Sun Feb 13 19:42:53 2011
@@ -4,6 +4,7 @@
"CodeGen"
"CodeGenCXX"
"CodeGenObjC"
+ "CodeGenOpenCL"
"Coverage"
"CXX"
"Driver"
Added: cfe/trunk/test/CodeGenOpenCL/kernel-metadata.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/kernel-metadata.cl?rev=125475&view=auto
==============================================================================
--- cfe/trunk/test/CodeGenOpenCL/kernel-metadata.cl (added)
+++ cfe/trunk/test/CodeGenOpenCL/kernel-metadata.cl Sun Feb 13 19:42:53 2011
@@ -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}
Added: cfe/trunk/test/Parser/opencl-kernel.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/Parser/opencl-kernel.cl?rev=125475&view=auto
==============================================================================
--- cfe/trunk/test/Parser/opencl-kernel.cl (added)
+++ cfe/trunk/test/Parser/opencl-kernel.cl Sun Feb 13 19:42:53 2011
@@ -0,0 +1,9 @@
+// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only
+
+__kernel void test()
+{
+}
+
+kernel void test1()
+{
+}
Added: cfe/trunk/test/Parser/opencl-pragma.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/Parser/opencl-pragma.cl?rev=125475&view=auto
==============================================================================
--- cfe/trunk/test/Parser/opencl-pragma.cl (added)
+++ cfe/trunk/test/Parser/opencl-pragma.cl Sun Feb 13 19:42:53 2011
@@ -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}}
More information about the cfe-commits
mailing list