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 b01a6b1..ad913a4 100644
--- a/include/clang/Basic/Attr.td
+++ b/include/clang/Basic/Attr.td
@@ -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">];
diff --git a/include/clang/Basic/DiagnosticSemaKinds.td b/include/clang/Basic/DiagnosticSemaKinds.td
index 188a5b5..3c5c9b3 100644
--- a/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/include/clang/Basic/DiagnosticSemaKinds.td
@@ -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<
diff --git a/include/clang/Sema/AttributeList.h b/include/clang/Sema/AttributeList.h
index bdb756b..bc1c2e2 100644
--- a/include/clang/Sema/AttributeList.h
+++ b/include/clang/Sema/AttributeList.h
@@ -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,
diff --git a/lib/Sema/AttributeList.cpp b/lib/Sema/AttributeList.cpp
index 4faa672..409e248 100644
--- a/lib/Sema/AttributeList.cpp
+++ b/lib/Sema/AttributeList.cpp
@@ -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);
}
diff --git a/lib/Sema/SemaDeclAttr.cpp b/lib/Sema/SemaDeclAttr.cpp
index 07db49e..fac47db 100644
--- a/lib/Sema/SemaDeclAttr.cpp
+++ b/lib/Sema/SemaDeclAttr.cpp
@@ -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.
diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt
index 42c63fc..09f6972 100644
--- a/test/CMakeLists.txt
+++ b/test/CMakeLists.txt
@@ -18,6 +18,7 @@
"Preprocessor"
"Rewriter"
"Sema"
+ "SemaCUDA"
"SemaCXX"
"SemaObjC"
"SemaObjCXX"
diff --git a/test/SemaCUDA/cuda.h b/test/SemaCUDA/cuda.h
new file mode 100644
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
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) {}