[AST] Enable expression of OpenCL language address spaces an attribute

[bader]

Summary

Enable a way to set OpenCL language address space using attributes
in addition to existing keywords.

Signed-off-by: Victor Lomuller [e-mail]

[Anastasia]

clang/include/clang/Basic/Attr.td

1123

We don't really use ocl prefix for opencl currently? Can it be opencl_private and etc?

[merge_guards_bot]

Build result: pass - 60453 tests passed, 0 failed and 726 were skipped.

Log files: console-log.txt, CMakeCache.txt

[bader]

Change attribute prefix: ocl_ -> opencl_.

Updated OpenCL spelling in the comments.

[bader]

clang/include/clang/Basic/Attr.td

1123

Sure. Done.

[merge_guards_bot]

Build result: pass - 60453 tests passed, 0 failed and 726 were skipped.

Log files: console-log.txt, CMakeCache.txt

[Anastasia]

LGTM! Thanks!

I presume OpenCL addr space logic won't apply in all cases in non-OpenCL compilations i.e. for example C++ because we enclose some of the logic under LangOpts checks.

[bader]

In D71005#1770338, @Anastasia wrote:

LGTM! Thanks!

I presume OpenCL addr space logic won't apply in all cases in non-OpenCL compilations i.e. for example C++ because we enclose some of the logic under LangOpts checks.

I think so too.

[aaron.ballman]

clang/include/clang/Basic/Attr.td

1123

This looks like we're adding two levels of namespace -- is there a reason this should be clang::opencl_private as opposed to opencl::private? Is there something clang-specific to these?

[bader]

clang/include/clang/Basic/Attr.td

1123

is there a reason this should be clang::opencl_private as opposed to opencl::private?

I'm okay with [[opencl::private]] as well. I have only one problem - currently OpenCL address spaces are exposed as keywords and using them in C++ breaks valid C++ code.

Is there something clang-specific to these?

I guess no (except that it's implemented in only Clang :-) ).

@aaron.ballman, I already committed this version, sorry about that. I'll open another review with attribute renaming.

[aaron.ballman]

clang/include/clang/Basic/Attr.td

1123

is there a reason this should be clang::opencl_private as opposed to opencl::private?

I'm okay with [[opencl::private]] as well. I have only one problem - currently OpenCL address spaces are exposed as keywords and using them in C++ breaks valid C++ code.

I'm not certain who controls the OpenCL spec, but this seems like it should be a decision that comes from there. Or is this functionality outside of the OpenCL spec?

In attributes, when an identifier can be interpreted as a keyword it is required to be interpreted as an identifier. We have the same issue with [[gnu::const]]. See http://eel.is/c++draft/dcl.attr.grammar#4.sentence-5

Is there something clang-specific to these?

I guess no (except that it's implemented in only Clang :-) ).

I mostly mean: will other implementations of OpenCL want to have the same functionality? Or is this something we expect other OpenCL implementations to largely ignore?

@aaron.ballman, I already committed this version, sorry about that. I'll open another review with attribute renaming.

The turnaround time between opening the review and closing it was pretty quick, but you had a reasonable LG so it's no big deal. We can continue chatting about it on this review until we decide that changes really are needed.

[bader]

is there a reason this should be clang::opencl_private as opposed to opencl::private?

I'm okay with [[opencl::private]] as well. I have only one problem - currently OpenCL address spaces are exposed as keywords and using them in C++ breaks valid C++ code.

I'm not certain who controls the OpenCL spec, but this seems like it should be a decision that comes from there. Or is this functionality outside of the OpenCL spec?

I guess I owe you some background here. OpenCL spec is controlled by Khronos organization, but Khronos provides only C-based language specification "OpenCL C" and I think this functionality is outside of the OpenCL spec.

Today Khronos organization rely on SPIR-V standard to enable high-level languages. This standard defines binary form of low-level intermediate language, which can be targeted by high-level language compiler like C++/Python/Java/Haskel/etc. High level level languages usually have separate working groups or standardization committees defining language features.

In addition to that SYCL working group within Khronos organization defines a C++-based abstraction interface to program accelerators like GPU, FPGA, DPS. SYCL interfaces do not expose any new C++ keywords or attributes and can be implemented as library using standard C++, but this implementation won't be able to offload execution of C++ code to an accelerator as standard C++ do not provide required facilities. To enable execution on accelerators, we implemented "SYCL device compiler" mode in clang and enhanced SYCL library with a few non-standard C++ extensions (e.g. attributes). Our implementation targets SPIR-V format for accelerated code and OpenCL API to communicate with accelerators. AFAIK, Codeplay team uses CUDA to enable execution of the SYCL code on NVIDIA GPUs.

In our implementation we are trying to re-use existing clang infrastructure developed for other GPU programming models OpenCL/OpenMP offload/CUDA/HIP, but due to design differences it can't always be re-used as is. For instance, if I understand it correctly, we can't use OpenCL address space attributes exposed as keywords as they might break valid C++ code - https://godbolt.org/z/fF5Ng5.
That's the basic motivation for this change.

In attributes, when an identifier can be interpreted as a keyword it is required to be interpreted as an identifier. We have the same issue with [[gnu::const]]. See http://eel.is/c++draft/dcl.attr.grammar#4.sentence-5

If I understand it correctly, it's the reason why const attribute has "GCC" spelling, instead of "Keyword" here https://github.com/llvm/llvm-project/blob/master/clang/include/clang/Basic/Attr.td#L940. Right?

will other implementations of OpenCL want to have the same functionality? Or is this something we expect other OpenCL implementations to largely ignore?

Exposing these attributes as non-keywords is useful for SYCL implementations. Actually this patch was developed by Codeplay team for ComputeCPP compiler - another SYCL implementation and "donated" by @Naghasan to our open source implementation to avoid implementation divergence.

To summarize: SYCL implementation need some instruments to express OpenCL address spaces in C++ code w/o breaking existing C++ code. We find that attributes fit well for this use case, but I'm open for alternative ideas.

[aaron.ballman]

In D71005#1770760, @bader wrote:

is there a reason this should be clang::opencl_private as opposed to opencl::private?

I'm okay with [[opencl::private]] as well. I have only one problem - currently OpenCL address spaces are exposed as keywords and using them in C++ breaks valid C++ code.

I'm not certain who controls the OpenCL spec, but this seems like it should be a decision that comes from there. Or is this functionality outside of the OpenCL spec?

I guess I owe you some background here. OpenCL spec is controlled by Khronos organization, but Khronos provides only C-based language specification "OpenCL C" and I think this functionality is outside of the OpenCL spec.

Today Khronos organization rely on SPIR-V standard to enable high-level languages. This standard defines binary form of low-level intermediate language, which can be targeted by high-level language compiler like C++/Python/Java/Haskel/etc. High level level languages usually have separate working groups or standardization committees defining language features.

In addition to that SYCL working group within Khronos organization defines a C++-based abstraction interface to program accelerators like GPU, FPGA, DPS. SYCL interfaces do not expose any new C++ keywords or attributes and can be implemented as library using standard C++, but this implementation won't be able to offload execution of C++ code to an accelerator as standard C++ do not provide required facilities. To enable execution on accelerators, we implemented "SYCL device compiler" mode in clang and enhanced SYCL library with a few non-standard C++ extensions (e.g. attributes). Our implementation targets SPIR-V format for accelerated code and OpenCL API to communicate with accelerators. AFAIK, Codeplay team uses CUDA to enable execution of the SYCL code on NVIDIA GPUs.

In our implementation we are trying to re-use existing clang infrastructure developed for other GPU programming models OpenCL/OpenMP offload/CUDA/HIP, but due to design differences it can't always be re-used as is. For instance, if I understand it correctly, we can't use OpenCL address space attributes exposed as keywords as they might break valid C++ code - https://godbolt.org/z/fF5Ng5.
That's the basic motivation for this change.

Thank you for the background! Now for some background of my own. :-)

Double square bracket-style attributes (which are supported in both C and C++) can only have a single level of vendor namespacing. So we can go with clang:: or opencl:: but cannot go with clang::opencl::. We control the clang vendor namespace for attributes, but we do not control the opencl vendor namespace (if one even exists). However, these attributes seem like they don't really "belong" to clang because they're OpenCL-specific moreso than clang-specific. That's why I was asking if it would make more sense to put them under an opencl namespace -- this allows Microsoft or GCC (or whatever other implementation) to implement the same semantics for their OpenCL implementations but without requiring them to introduce our vendor namespace into their products. However, I wasn't certain what entity could make decisions about that vendor namespace. Based on what you wrote, it sounds like the Khronos group is the most likely one to want to control that namespace, but I'm not certain.

To be clear, there's nothing invalid or incorrect about clang::opencl_whatever, it just looks strange because that's basically adding a second-level namespace.

In attributes, when an identifier can be interpreted as a keyword it is required to be interpreted as an identifier. We have the same issue with [[gnu::const]]. See http://eel.is/c++draft/dcl.attr.grammar#4.sentence-5

If I understand it correctly, it's the reason why const attribute has "GCC" spelling, instead of "Keyword" here https://github.com/llvm/llvm-project/blob/master/clang/include/clang/Basic/Attr.td#L940. Right?

I may not have been sufficiently clear. I was talking about attributes within double square brackets, like [[gnu::const]] -- that const is interpreted as an identifier rather than a keyword, making this a legal attribute in C and C++. So there's nothing wrong with [[opencl::private]] as an attribute, despite private being a keyword. I think you're talking about turning the keyword-like attribute __private into private, then you're correct that we do not have a way to support something like that.

will other implementations of OpenCL want to have the same functionality? Or is this something we expect other OpenCL implementations to largely ignore?

Exposing these attributes as non-keywords is useful for SYCL implementations. Actually this patch was developed by Codeplay team for ComputeCPP compiler - another SYCL implementation and "donated" by @Naghasan to our open source implementation to avoid implementation divergence.

To summarize: SYCL implementation need some instruments to express OpenCL address spaces in C++ code w/o breaking existing C++ code. We find that attributes fit well for this use case, but I'm open for alternative ideas.

I think attributes are a great approach to solving the problem. I'm only pushing back on the names of the attributes and whether we think they should be in the clang vendor namespace or an opencl one. Now that C2x has support for attributes, I would think that Khronos could be approached to see if they would like to see [[opencl::private]] as an attribute in C that could then also be used with the same semantics in a C++ program. WDYT? If they wanted to "own" the opencl vendor namespace, that would make me more comfortable adding a new attribute under that namespace. If they don't want to own that namespace, or if they don't want to use the attribute syntax in C at all, that's also fine -- we can go with what's already committed in the clang vendor namespace.

[saugustine]

When I try to run tablegen on this file for clang/docs (as described in clang/docs/InternalsManual.rst), I get an error, introduced by this change:

$ ./bin/clang-tblgen -gen-attr-docs -I ../llvm-project/clang/include ../llvm-project/clang/include/clang/Basic/Attr.td -o /tmp/AttributeReference.rst

../llvm-project/clang/include/clang/Basic/Attr.td:1137:1: error: This attribute requires a heading to be specified
def OpenCLConstantAddressSpace : TypeAttr {
^

I'm at a loss as to what is wrong, but I would like to revert this change until it is fixed.

[aaron.ballman]

In D71005#1771513, @saugustine wrote:

When I try to run tablegen on this file for clang/docs (as described in clang/docs/InternalsManual.rst), I get an error, introduced by this change:

$ ./bin/clang-tblgen -gen-attr-docs -I ../llvm-project/clang/include ../llvm-project/clang/include/clang/Basic/Attr.td -o /tmp/AttributeReference.rst

../llvm-project/clang/include/clang/Basic/Attr.td:1137:1: error: This attribute requires a heading to be specified
def OpenCLConstantAddressSpace : TypeAttr {
^

I'm at a loss as to what is wrong, but I would like to revert this change until it is fixed.

This should be fixed in f5193d87feaedb411255e92979abd6b62522bc38 now, thank you for pointing it out!

[bader]

@aaron.ballman, thank you for fixing the problem with documentation generation.

With regards to attributes naming question, let me check if I get it right.

1. You suggest checking with Khronos if they want to adopt [[opencl::private]] attribute names in OpenCL kernel language standard. Right? @Anastasia, can you help with this?
2. If Khronos leaves this decision to implementers, are current names (e.g. [[clang::opencl_private]]) look okay or we should consider other options?

[aaron.ballman]

In D71005#1773042, @bader wrote:

@aaron.ballman, thank you for fixing the problem with documentation generation.

No problem!

With regards to attributes naming question, let me check if I get it right.

1. You suggest checking with Khronos if they want to adopt [[opencl::private]] attribute names in OpenCL kernel language standard. Right? @Anastasia, can you help with this?

That's correct.

2. If Khronos leaves this decision to implementers, are current names (e.g. [[clang::opencl_private]]) look okay or we should consider other options?

If they want to leave it to implementers, I think the names you have now are fine.

clang/include/clang/Basic/Attr.td

let Accessors = [Accessor<"isReadOnly", [Keyword<"__read_only">,
                    Accessor<"isReadWrite", [Keyword<"__read_write">,
                                             Keyword<"read_write">]>,
                    Accessor<"isWriteOnly", [Keyword<"__write_only">,
                                             Keyword<"write_only">]>];
   let Documentation = [OpenCLAccessDocs];
 }
 
 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];
 }
 
 def OpenCLNoSVM : Attr {
   let Spellings = [GNU<"nosvm">];
   let Subjects = SubjectList<[Var]>;
   let Documentation = [OpenCLNoSVMDocs];
   let LangOpts = [OpenCL];

clang/lib/Sema/SemaType.cpp

if (State.getDeclarator().isDeclarationOfFunction()) {
     CurType = State.getAttributedType(
         createSimpleAttr<LifetimeBoundAttr>(State.getSema().Context, Attr),
         CurType, CurType);
   } else {
     Attr.diagnoseAppertainsTo(State.getSema(), nullptr);
   }
 }
 
+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,
                              ParsedAttributesView &attrs) {
   // Scan through and apply attributes to this type where it makes sense.  Some
   // attributes (such as __address_space__, __vector_size__, etc) apply to the
   // type, but others can be present in the type specifiers even though they
   // apply to the decl.  Here we apply type attributes and ignore the rest.
if (attr.isCXX11Attribute()) {
           state.getSema().Diag(attr.getLoc(),
                                IsTypeAttr
                                    ? diag::warn_gcc_ignores_type_attr
                                    : diag::warn_cxx11_gnu_attribute_on_type)
               << attr;
           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;
       }
     }
 
     // If this is an attribute we can handle, do so now,
     // otherwise, add it to the FnAttrs list for rechaining.
     switch (attr.getKind()) {
     default:

clang/test/AST/language_address_space_attribute.cpp

+// 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

 }
 
 void func_with_array_param(const unsigned data[16]);
 
 __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__
+}