[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@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@intel.com>

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];
 }
 

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;
       }
     }

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;
+}

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__
+}