Author: Victor Lomuller Date: 2019-12-05T12:24:06+03:00 New Revision: 11a9bae8f66986751078501988b4414f24dbe37e
URL: https://github.com/llvm/llvm-project/commit/11a9bae8f66986751078501988b4414f24dbe37e DIFF: https://github.com/llvm/llvm-project/commit/11a9bae8f66986751078501988b4414f24dbe37e.diff LOG: [AST] Enable expression of OpenCL language address spaces an attribute Summary: Enable a way to set OpenCL language address space using attributes in addition to existing keywords. Signed-off-by: Victor Lomuller vic...@codeplay.com Reviewers: aaron.ballman, Anastasia Subscribers: yaxunl, ebevhan, cfe-commits, Naghasan Tags: #clang Differential Revision: https://reviews.llvm.org/D71005 Signed-off-by: Alexey Bader <alexey.ba...@intel.com> Added: clang/test/AST/language_address_space_attribute.cpp Modified: clang/include/clang/Basic/Attr.td clang/lib/Sema/SemaType.cpp clang/test/SemaOpenCL/address-spaces.cl Removed: ################################################################################ diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index 4ea1c9f58beb..9ca4be0e07c8 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 52a0581643bc..1375ccbabc50 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 000000000000..7c6bdca06c06 --- /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 09a6dd0ba53f..a28069470177 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__ +} _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
