diff options
Diffstat (limited to 'clang')
| -rw-r--r-- | clang/include/clang/Basic/Attr.td | 10 | ||||
| -rw-r--r-- | clang/lib/Sema/SemaType.cpp | 16 | ||||
| -rw-r--r-- | clang/test/AST/language_address_space_attribute.cpp | 36 | ||||
| -rw-r--r-- | clang/test/SemaOpenCL/address-spaces.cl | 16 |
4 files changed, 70 insertions, 8 deletions
diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index 4ea1c9f58be..9ca4be0e07c 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -1120,27 +1120,27 @@ def OpenCLAccess : Attr { } def OpenCLPrivateAddressSpace : TypeAttr { - let Spellings = [Keyword<"__private">, Keyword<"private">]; + let Spellings = [Keyword<"__private">, Keyword<"private">, Clang<"opencl_private">]; let Documentation = [OpenCLAddressSpacePrivateDocs]; } def OpenCLGlobalAddressSpace : TypeAttr { - let Spellings = [Keyword<"__global">, Keyword<"global">]; + let Spellings = [Keyword<"__global">, Keyword<"global">, Clang<"opencl_global">]; let Documentation = [OpenCLAddressSpaceGlobalDocs]; } def OpenCLLocalAddressSpace : TypeAttr { - let Spellings = [Keyword<"__local">, Keyword<"local">]; + let Spellings = [Keyword<"__local">, Keyword<"local">, Clang<"opencl_local">]; let Documentation = [OpenCLAddressSpaceLocalDocs]; } def OpenCLConstantAddressSpace : TypeAttr { - let Spellings = [Keyword<"__constant">, Keyword<"constant">]; + let Spellings = [Keyword<"__constant">, Keyword<"constant">, Clang<"opencl_constant">]; let Documentation = [OpenCLAddressSpaceConstantDocs]; } def OpenCLGenericAddressSpace : TypeAttr { - let Spellings = [Keyword<"__generic">, Keyword<"generic">]; + let Spellings = [Keyword<"__generic">, Keyword<"generic">, Clang<"opencl_generic">]; let Documentation = [OpenCLAddressSpaceGenericDocs]; } diff --git a/clang/lib/Sema/SemaType.cpp b/clang/lib/Sema/SemaType.cpp index 52a0581643b..1375ccbabc5 100644 --- a/clang/lib/Sema/SemaType.cpp +++ b/clang/lib/Sema/SemaType.cpp @@ -7407,6 +7407,16 @@ static void HandleLifetimeBoundAttr(TypeProcessingState &State, } } +static bool isAddressSpaceKind(const ParsedAttr &attr) { + auto attrKind = attr.getKind(); + + return attrKind == ParsedAttr::AT_AddressSpace || + attrKind == ParsedAttr::AT_OpenCLPrivateAddressSpace || + attrKind == ParsedAttr::AT_OpenCLGlobalAddressSpace || + attrKind == ParsedAttr::AT_OpenCLLocalAddressSpace || + attrKind == ParsedAttr::AT_OpenCLConstantAddressSpace || + attrKind == ParsedAttr::AT_OpenCLGenericAddressSpace; +} static void processTypeAttrs(TypeProcessingState &state, QualType &type, TypeAttrLocation TAL, @@ -7445,11 +7455,11 @@ static void processTypeAttrs(TypeProcessingState &state, QualType &type, if (!IsTypeAttr) continue; } - } else if (TAL != TAL_DeclChunk && - attr.getKind() != ParsedAttr::AT_AddressSpace) { + } else if (TAL != TAL_DeclChunk && !isAddressSpaceKind(attr)) { // Otherwise, only consider type processing for a C++11 attribute if // it's actually been applied to a type. - // We also allow C++11 address_space attributes to pass through. + // We also allow C++11 address_space and + // OpenCL language address space attributes to pass through. continue; } } diff --git a/clang/test/AST/language_address_space_attribute.cpp b/clang/test/AST/language_address_space_attribute.cpp new file mode 100644 index 00000000000..7c6bdca06c0 --- /dev/null +++ b/clang/test/AST/language_address_space_attribute.cpp @@ -0,0 +1,36 @@ +// RUN: %clang_cc1 %s -ast-dump | FileCheck %s + +// Verify that the language address space attribute is +// understood correctly by clang. + +void langas() { + // CHECK: VarDecl {{.*}} x_global '__global int *' + __attribute__((opencl_global)) int *x_global; + + // CHECK: VarDecl {{.*}} z_global '__global int *' + [[clang::opencl_global]] int *z_global; + + // CHECK: VarDecl {{.*}} x_local '__local int *' + __attribute__((opencl_local)) int *x_local; + + // CHECK: VarDecl {{.*}} z_local '__local int *' + [[clang::opencl_local]] int *z_local; + + // CHECK: VarDecl {{.*}} x_constant '__constant int *' + __attribute__((opencl_constant)) int *x_constant; + + // CHECK: VarDecl {{.*}} z_constant '__constant int *' + [[clang::opencl_constant]] int *z_constant; + + // CHECK: VarDecl {{.*}} x_private 'int *' + __attribute__((opencl_private)) int *x_private; + + // CHECK: VarDecl {{.*}} z_private 'int *' + [[clang::opencl_private]] int *z_private; + + // CHECK: VarDecl {{.*}} x_generic '__generic int *' + __attribute__((opencl_generic)) int *x_generic; + + // CHECK: VarDecl {{.*}} z_generic '__generic int *' + [[clang::opencl_generic]] int *z_generic; +} diff --git a/clang/test/SemaOpenCL/address-spaces.cl b/clang/test/SemaOpenCL/address-spaces.cl index 09a6dd0ba53..a2806947017 100644 --- a/clang/test/SemaOpenCL/address-spaces.cl +++ b/clang/test/SemaOpenCL/address-spaces.cl @@ -248,3 +248,19 @@ __kernel void k() { unsigned data[16]; func_with_array_param(data); } + +void func_multiple_addr2(void) { + typedef __private int private_int_t; + __private __attribute__((opencl_global)) int var1; // expected-error {{multiple address spaces specified for type}} + __private __attribute__((opencl_global)) int *var2; // expected-error {{multiple address spaces specified for type}} + __attribute__((opencl_global)) private_int_t var3; // expected-error {{multiple address spaces specified for type}} + __attribute__((opencl_global)) private_int_t *var4; // expected-error {{multiple address spaces specified for type}} + __attribute__((opencl_private)) private_int_t var5; // expected-warning {{multiple identical address spaces specified for type}} + __attribute__((opencl_private)) private_int_t *var6; // expected-warning {{multiple identical address spaces specified for type}} +#if __OPENCL_CPP_VERSION__ + [[clang::opencl_private]] __global int var7; // expected-error {{multiple address spaces specified for type}} + [[clang::opencl_private]] __global int *var8; // expected-error {{multiple address spaces specified for type}} + [[clang::opencl_private]] private_int_t var9; // expected-warning {{multiple identical address spaces specified for type}} + [[clang::opencl_private]] private_int_t *var10; // expected-warning {{multiple identical address spaces specified for type}} +#endif // !__OPENCL_CPP_VERSION__ +} |
