aboutsummaryrefslogtreecommitdiff
diff options
context:
space:
mode:
Diffstat
-rw-r--r--include/clang/Basic/Attr.td20
-rw-r--r--include/clang/Basic/DiagnosticSemaKinds.td4
-rw-r--r--include/clang/Sema/AttributeList.h5
-rw-r--r--lib/Sema/AttributeList.cpp5
-rw-r--r--lib/Sema/SemaDeclAttr.cpp105
-rw-r--r--test/CMakeLists.txt1
-rw-r--r--test/SemaCUDA/cuda.h7
-rw-r--r--test/SemaCUDA/qualifiers.cu5
8 files changed, 150 insertions, 2 deletions
diff --git a/include/clang/Basic/Attr.td b/include/clang/Basic/Attr.td
index b01a6b1a7d..ad913a4901 100644
--- 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 188a5b5909..3c5c9b32e5 100644
--- 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 bdb756b329..bc1c2e29fb 100644
--- 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 4faa67223c..409e2488bd 100644
--- 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 07db49eb26..fac47db7f9 100644
--- 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 42c63fccf4..09f69720e0 100644
--- 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
index 0000000000..c503747820
--- /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 0000000000..8d5b759a6d
--- /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) {}
generated by cgit v1.2.3-18-g5258 (git 2.32.0) at 2025-07-24 21:07:11 +0000