[clang] 11a9bae - [AST] Enable expression of OpenCL language address spaces an attribute
Alexey Bader via cfe-commits
cfe-commits at lists.llvm.org
Thu Dec 5 04:48:09 PST 2019
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 victor at 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.bader at 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__
+}
More information about the cfe-commits
mailing list