[cfe-commits] r120545 - in /cfe/trunk: include/clang/Basic/Attr.td include/clang/Basic/DiagnosticSemaKinds.td include/clang/Sema/AttributeList.h lib/Sema/AttributeList.cpp lib/Sema/SemaDeclAttr.cpp test/CMakeLists.txt test/SemaCUDA/ test/SemaCUDA/cuda.h test/SemaCUDA/qualifiers.cu
Peter Collingbourne
peter at pcc.me.uk
Tue Nov 30 19:15:32 PST 2010
Author: pcc
Date: Tue Nov 30 21:15:31 2010
New Revision: 120545
URL: http://llvm.org/viewvc/llvm-project?rev=120545&view=rev
Log:
Basic, Sema: add support for CUDA location attributes
Added:
cfe/trunk/test/SemaCUDA/
cfe/trunk/test/SemaCUDA/cuda.h
cfe/trunk/test/SemaCUDA/qualifiers.cu
Modified:
cfe/trunk/include/clang/Basic/Attr.td
cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
cfe/trunk/include/clang/Sema/AttributeList.h
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=120545&r1=120544&r2=120545&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/Attr.td (original)
+++ cfe/trunk/include/clang/Basic/Attr.td Tue Nov 30 21:15:31 2010
@@ -170,6 +170,26 @@
let Args = [IntArgument<"Priority">];
}
+def CUDAConstant : Attr {
+ let Spellings = ["constant"];
+}
+
+def CUDADevice : Attr {
+ let Spellings = ["device"];
+}
+
+def CUDAGlobal : Attr {
+ let Spellings = ["global"];
+}
+
+def CUDAHost : Attr {
+ let Spellings = ["host"];
+}
+
+def CUDAShared : Attr {
+ let Spellings = ["shared"];
+}
+
def Deprecated : Attr {
let Spellings = ["deprecated"];
let Args = [StringArgument<"Message">];
Modified: cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td?rev=120545&r1=120544&r2=120545&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td (original)
+++ cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td Tue Nov 30 21:15:31 2010
@@ -1006,13 +1006,13 @@
"variable and function|function or method|parameter|"
"parameter or Objective-C method |function, method or block|"
"virtual method or class|function, method, or parameter|class|virtual method"
- "|member}1 types">;
+ "|member|variable}1 types">;
def err_attribute_wrong_decl_type : Error<
"%0 attribute only applies to %select{function|union|"
"variable and function|function or method|parameter|"
"parameter or Objective-C method |function, method or block|"
"virtual method or class|function, method, or parameter|class|virtual method"
- "|member}1 types">;
+ "|member|variable}1 types">;
def warn_function_attribute_wrong_type : Warning<
"%0 only applies to function types; type here is %1">;
def warn_gnu_inline_attribute_requires_inline : Warning<
Modified: cfe/trunk/include/clang/Sema/AttributeList.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Sema/AttributeList.h?rev=120545&r1=120544&r2=120545&view=diff
==============================================================================
--- cfe/trunk/include/clang/Sema/AttributeList.h (original)
+++ cfe/trunk/include/clang/Sema/AttributeList.h Tue Nov 30 21:15:31 2010
@@ -92,9 +92,11 @@
AT_cdecl,
AT_cleanup,
AT_const,
+ AT_constant,
AT_constructor,
AT_deprecated,
AT_destructor,
+ AT_device,
AT_dllexport,
AT_dllimport,
AT_ext_vector_type,
@@ -102,8 +104,10 @@
AT_final,
AT_format,
AT_format_arg,
+ AT_global,
AT_gnu_inline,
AT_hiding,
+ AT_host,
AT_malloc,
AT_may_alias,
AT_mode,
@@ -134,6 +138,7 @@
AT_regparm,
AT_section,
AT_sentinel,
+ AT_shared,
AT_stdcall,
AT_thiscall,
AT_transparent_union,
Modified: cfe/trunk/lib/Sema/AttributeList.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/AttributeList.cpp?rev=120545&r1=120544&r2=120545&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/AttributeList.cpp (original)
+++ cfe/trunk/lib/Sema/AttributeList.cpp Tue Nov 30 21:15:31 2010
@@ -125,5 +125,10 @@
.Case("__fastcall", AT_fastcall)
.Case("__thiscall", AT_thiscall)
.Case("__pascal", AT_pascal)
+ .Case("constant", AT_constant)
+ .Case("device", AT_device)
+ .Case("global", AT_global)
+ .Case("host", AT_host)
+ .Case("shared", AT_shared)
.Default(UnknownAttribute);
}
Modified: cfe/trunk/lib/Sema/SemaDeclAttr.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaDeclAttr.cpp?rev=120545&r1=120544&r2=120545&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaDeclAttr.cpp (original)
+++ cfe/trunk/lib/Sema/SemaDeclAttr.cpp Tue Nov 30 21:15:31 2010
@@ -2078,6 +2078,106 @@
d->addAttr(::new (S.Context) NoInstrumentFunctionAttr(Attr.getLoc(), S.Context));
}
+static void HandleConstantAttr(Decl *d, const AttributeList &Attr, Sema &S) {
+ if (S.LangOpts.CUDA) {
+ // check the attribute arguments.
+ if (Attr.getNumArgs() != 0) {
+ S.Diag(Attr.getLoc(), diag::err_attribute_wrong_number_arguments) << 0;
+ return;
+ }
+
+ if (!isa<VarDecl>(d)) {
+ S.Diag(Attr.getLoc(), diag::warn_attribute_wrong_decl_type)
+ << Attr.getName() << 12 /*variable*/;
+ return;
+ }
+
+ d->addAttr(::new (S.Context) CUDAConstantAttr(Attr.getLoc(), S.Context));
+ } else {
+ S.Diag(Attr.getLoc(), diag::warn_attribute_ignored) << "constant";
+ }
+}
+
+static void HandleDeviceAttr(Decl *d, const AttributeList &Attr, Sema &S) {
+ if (S.LangOpts.CUDA) {
+ // check the attribute arguments.
+ if (Attr.getNumArgs() != 0) {
+ S.Diag(Attr.getLoc(), diag::err_attribute_wrong_number_arguments) << 0;
+ return;
+ }
+
+ if (!isa<FunctionDecl>(d) && !isa<VarDecl>(d)) {
+ S.Diag(Attr.getLoc(), diag::warn_attribute_wrong_decl_type)
+ << Attr.getName() << 2 /*variable and function*/;
+ return;
+ }
+
+ d->addAttr(::new (S.Context) CUDADeviceAttr(Attr.getLoc(), S.Context));
+ } else {
+ S.Diag(Attr.getLoc(), diag::warn_attribute_ignored) << "device";
+ }
+}
+
+static void HandleGlobalAttr(Decl *d, const AttributeList &Attr, Sema &S) {
+ if (S.LangOpts.CUDA) {
+ // check the attribute arguments.
+ if (Attr.getNumArgs() != 0) {
+ S.Diag(Attr.getLoc(), diag::err_attribute_wrong_number_arguments) << 0;
+ return;
+ }
+
+ if (!isa<FunctionDecl>(d)) {
+ S.Diag(Attr.getLoc(), diag::warn_attribute_wrong_decl_type)
+ << Attr.getName() << 0 /*function*/;
+ return;
+ }
+
+ d->addAttr(::new (S.Context) CUDAGlobalAttr(Attr.getLoc(), S.Context));
+ } else {
+ S.Diag(Attr.getLoc(), diag::warn_attribute_ignored) << "global";
+ }
+}
+
+static void HandleHostAttr(Decl *d, const AttributeList &Attr, Sema &S) {
+ if (S.LangOpts.CUDA) {
+ // check the attribute arguments.
+ if (Attr.getNumArgs() != 0) {
+ S.Diag(Attr.getLoc(), diag::err_attribute_wrong_number_arguments) << 0;
+ return;
+ }
+
+ if (!isa<FunctionDecl>(d)) {
+ S.Diag(Attr.getLoc(), diag::warn_attribute_wrong_decl_type)
+ << Attr.getName() << 0 /*function*/;
+ return;
+ }
+
+ d->addAttr(::new (S.Context) CUDAHostAttr(Attr.getLoc(), S.Context));
+ } else {
+ S.Diag(Attr.getLoc(), diag::warn_attribute_ignored) << "host";
+ }
+}
+
+static void HandleSharedAttr(Decl *d, const AttributeList &Attr, Sema &S) {
+ if (S.LangOpts.CUDA) {
+ // check the attribute arguments.
+ if (Attr.getNumArgs() != 0) {
+ S.Diag(Attr.getLoc(), diag::err_attribute_wrong_number_arguments) << 0;
+ return;
+ }
+
+ if (!isa<VarDecl>(d)) {
+ S.Diag(Attr.getLoc(), diag::warn_attribute_wrong_decl_type)
+ << Attr.getName() << 12 /*variable*/;
+ return;
+ }
+
+ d->addAttr(::new (S.Context) CUDASharedAttr(Attr.getLoc(), S.Context));
+ } else {
+ S.Diag(Attr.getLoc(), diag::warn_attribute_ignored) << "shared";
+ }
+}
+
static void HandleGNUInlineAttr(Decl *d, const AttributeList &Attr, Sema &S) {
// check the attribute arguments.
if (Attr.getNumArgs() != 0) {
@@ -2358,17 +2458,21 @@
case AttributeList::AT_base_check: HandleBaseCheckAttr (D, Attr, S); break;
case AttributeList::AT_carries_dependency:
HandleDependencyAttr (D, Attr, S); break;
+ case AttributeList::AT_constant: HandleConstantAttr (D, Attr, S); break;
case AttributeList::AT_constructor: HandleConstructorAttr (D, Attr, S); break;
case AttributeList::AT_deprecated: HandleDeprecatedAttr (D, Attr, S); break;
case AttributeList::AT_destructor: HandleDestructorAttr (D, Attr, S); break;
+ case AttributeList::AT_device: HandleDeviceAttr (D, Attr, S); break;
case AttributeList::AT_ext_vector_type:
HandleExtVectorTypeAttr(scope, D, Attr, S);
break;
case AttributeList::AT_final: HandleFinalAttr (D, Attr, S); break;
case AttributeList::AT_format: HandleFormatAttr (D, Attr, S); break;
case AttributeList::AT_format_arg: HandleFormatArgAttr (D, Attr, S); break;
+ case AttributeList::AT_global: HandleGlobalAttr (D, Attr, S); break;
case AttributeList::AT_gnu_inline: HandleGNUInlineAttr (D, Attr, S); break;
case AttributeList::AT_hiding: HandleHidingAttr (D, Attr, S); break;
+ case AttributeList::AT_host: HandleHostAttr (D, Attr, S); break;
case AttributeList::AT_mode: HandleModeAttr (D, Attr, S); break;
case AttributeList::AT_malloc: HandleMallocAttr (D, Attr, S); break;
case AttributeList::AT_may_alias: HandleMayAliasAttr (D, Attr, S); break;
@@ -2381,6 +2485,7 @@
case AttributeList::AT_noreturn: HandleNoReturnAttr (D, Attr, S); break;
case AttributeList::AT_nothrow: HandleNothrowAttr (D, Attr, S); break;
case AttributeList::AT_override: HandleOverrideAttr (D, Attr, S); break;
+ case AttributeList::AT_shared: HandleSharedAttr (D, Attr, S); break;
case AttributeList::AT_vecreturn: HandleVecReturnAttr (D, Attr, S); break;
// Checker-specific.
Modified: cfe/trunk/test/CMakeLists.txt
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CMakeLists.txt?rev=120545&r1=120544&r2=120545&view=diff
==============================================================================
--- cfe/trunk/test/CMakeLists.txt (original)
+++ cfe/trunk/test/CMakeLists.txt Tue Nov 30 21:15:31 2010
@@ -18,6 +18,7 @@
"Preprocessor"
"Rewriter"
"Sema"
+ "SemaCUDA"
"SemaCXX"
"SemaObjC"
"SemaObjCXX"
Added: cfe/trunk/test/SemaCUDA/cuda.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/cuda.h?rev=120545&view=auto
==============================================================================
--- cfe/trunk/test/SemaCUDA/cuda.h (added)
+++ cfe/trunk/test/SemaCUDA/cuda.h Tue Nov 30 21:15:31 2010
@@ -0,0 +1,7 @@
+/* Minimal declarations for CUDA support. Testing purposes only. */
+
+#define __constant__ __attribute__((constant))
+#define __device__ __attribute__((device))
+#define __global__ __attribute__((global))
+#define __host__ __attribute__((host))
+#define __shared__ __attribute__((shared))
Added: cfe/trunk/test/SemaCUDA/qualifiers.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/qualifiers.cu?rev=120545&view=auto
==============================================================================
--- cfe/trunk/test/SemaCUDA/qualifiers.cu (added)
+++ cfe/trunk/test/SemaCUDA/qualifiers.cu Tue Nov 30 21:15:31 2010
@@ -0,0 +1,5 @@
+// RUN: %clang_cc1 -fsyntax-only -verify %s
+
+#include "cuda.h"
+
+__global__ void g1(int x) {}
More information about the cfe-commits
mailing list