Basic, Sema: add support for CUDA location attributes

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

diff --git a/include/clang/Basic/Attr.td b/include/clang/Basic/Attr.td
index b01a6b1a7d62f72488385ea203c150b06ab5f597..ad913a49011d306d8c1e630a9886cdfa0c4f52e9 100644 (file)
--- a/include/clang/Basic/Attr.td
+++ b/include/clang/Basic/Attr.td
@@ -170,6 +170,26 @@ def Constructor : Attr {
   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">];

diff --git a/include/clang/Basic/DiagnosticSemaKinds.td b/include/clang/Basic/DiagnosticSemaKinds.td
index 188a5b59090948cb3380d5f0e46dbfd65b52dc14..3c5c9b32e50e1c0e86e039888e187654d2545c22 100644 (file)
--- a/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/include/clang/Basic/DiagnosticSemaKinds.td
@@ -1006,13 +1006,13 @@ def warn_attribute_wrong_decl_type : Warning<
   "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<

diff --git a/include/clang/Sema/AttributeList.h b/include/clang/Sema/AttributeList.h
index bdb756b329c220aa5bef3ef35f6f693b93bd5923..bc1c2e29fbf83c71871c05b84f6fc26b165bd8d0 100644 (file)
--- a/include/clang/Sema/AttributeList.h
+++ b/include/clang/Sema/AttributeList.h
@@ -92,9 +92,11 @@ public:
     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 @@ public:
     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 @@ public:
     AT_regparm,
     AT_section,
     AT_sentinel,
+    AT_shared,
     AT_stdcall,
     AT_thiscall,
     AT_transparent_union,

diff --git a/lib/Sema/AttributeList.cpp b/lib/Sema/AttributeList.cpp
index 4faa67223c36253a8cbbb7ac57a0d24648b10c8e..409e2488bde375236c192c5963372f4e679d3790 100644 (file)
--- a/lib/Sema/AttributeList.cpp
+++ b/lib/Sema/AttributeList.cpp
@@ -125,5 +125,10 @@ AttributeList::Kind AttributeList::getKind(const IdentifierInfo *Name) {
     .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);
 }

diff --git a/lib/Sema/SemaDeclAttr.cpp b/lib/Sema/SemaDeclAttr.cpp
index 07db49eb262994156bd69de93bb44a69e87af429..fac47db7f936de7d50f53552b5c55e967c3da288 100644 (file)
--- a/lib/Sema/SemaDeclAttr.cpp
+++ b/lib/Sema/SemaDeclAttr.cpp
@@ -2078,6 +2078,106 @@ static void HandleNoInstrumentFunctionAttr(Decl *d, const AttributeList &Attr,
   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 @@ static void ProcessDeclAttribute(Scope *scope, Decl *D,
   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 @@ static void ProcessDeclAttribute(Scope *scope, Decl *D,
   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.

diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt
index 42c63fccf4f513ab905a203e94b4d074fb861056..09f69720e0ec9145f56506f659032bd92ee00250 100644 (file)
--- a/test/CMakeLists.txt
+++ b/test/CMakeLists.txt
@@ -18,6 +18,7 @@ set(CLANG_TEST_DIRECTORIES
   "Preprocessor"
   "Rewriter"
   "Sema"
+  "SemaCUDA"
   "SemaCXX"
   "SemaObjC"
   "SemaObjCXX"

diff --git a/test/SemaCUDA/cuda.h b/test/SemaCUDA/cuda.h
new file mode 100644 (file)
index 0000000..c503747
--- /dev/null
+++ b/test/SemaCUDA/cuda.h
@@ -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))

diff --git a/test/SemaCUDA/qualifiers.cu b/test/SemaCUDA/qualifiers.cu
new file mode 100644 (file)
index 0000000..8d5b759
--- /dev/null
+++ b/test/SemaCUDA/qualifiers.cu
@@ -0,0 +1,5 @@
+// RUN: %clang_cc1 -fsyntax-only -verify %s
+
+#include "cuda.h"
+
+__global__ void g1(int x) {}