From: Peter Collingbourne Date: Wed, 1 Dec 2010 03:15:31 +0000 (+0000) Subject: Basic, Sema: add support for CUDA location attributes X-Git-Url: https://granicus.if.org/sourcecode?a=commitdiff_plain;h=ced7671c18e115ac3c3f54abfaaafcc6d33edc4c;p=clang 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 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(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(d) && !isa(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(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(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(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) {}