[PATCH] D60455: [SYCL] Add support for SYCL device attributes

[Aaron Ballman via Phabricator via cfe-commits]

aaron.ballman requested changes to this revision.
aaron.ballman added a comment.
This revision now requires changes to proceed.

In D60455#1464324 <https://reviews.llvm.org/D60455#1464324>, @Fznamznon wrote:

> Applied comments from @aaron.ballman and @keryell
>
> - Introduced a C++11 and C2x style spelling in the clang namespace. I didn't find path to add two namespaces to attribute (like [[clang::sycl::device]]) so [[clang::sycl_device]] spelling is added.
> - Since both attributes have spellings and possible can be used as some "standard" outlining in Clang/LLVM I added documetation.
> - Added more test cases.
> - As @bader mentioned sycl_device can be used to mark functions, which are called from the different translation units so I added simple handling of this attribute in Sema.


I'm confused -- I thought @bader also said "...SYCL is not supposed to expose any non-standard extensions to a user." -- these attributes are not standards based (WG21 and WG14 have nothing to say about them), so are these attributes considered "non-standard extensions" or not?

I'm still wondering what the actual semantics are for the attribute. Right now, these are being parsed and ignored -- are there future plans here?



================
Comment at: clang/include/clang/Basic/Attr.td:1000
+def SYCLDevice : InheritableAttr {
+  let Spellings = [GNU<"sycl_device">];
+  let Subjects = SubjectList<[Function, Var]>;
----------------
aaron.ballman wrote:
> keryell wrote:
> > Fznamznon wrote:
> > > aaron.ballman wrote:
> > > > Is there a reason to not also introduce a C++11 and C2x style spelling in the `clang` namespace? e.g., `[[clang::sycl_device]]`
> > > I don't think that it makes sense because these attributes not for public consumption. These attributes is needed to separate code which is supposed to be offloaded from regular host code. I think SYCLDevice attribute actually doesn't need a spelling because it will be added only implicitly by compiler.
> > 
> > If we go towards this direction, `[[clang::sycl::device]]` or `[[clang::sycl::kernel]]` look more compatible with the concept of name space.
> > While not a public interface, if we have a kind of "standard" outlining in Clang/LLVM, some people might want to use it in some other contexts too.
> If these are only being added implicitly by the compiler, then they should not be given any `Spelling`. See `AlignMac68k` for an example.
I'm still confused -- are these created implicitly or are they spelled out by the user explicitly? Right now, it looks like they're spelled out explicitly, but I was under the impression they are only intended to be created implicitly by the compiler.

If they are expected to be explicitly specified by the user, the spelling should be using `Clang<>` instead of using `GNU<>`, `C2x<>`, and `CXX11<>` explicitly.

> If we go towards this direction, [[clang::sycl::device]] or [[clang::sycl::kernel]] look more compatible with the concept of name space.

Attribute namespaces do not work that way. There is the vendor namespace and then the attribute name.


================
Comment at: clang/include/clang/Basic/AttrDocs.td:259
+  let Content = [{
+The sycl_device attribute specifies function which is supposed to be compiled
+for the device and cannot be directly called by the host. Here is code example
----------------
specifies function which is -> specifies that a function is

Also, please put backticks around `sycl_device`.


================
Comment at: clang/include/clang/Basic/AttrDocs.td:260
+The sycl_device attribute specifies function which is supposed to be compiled
+for the device and cannot be directly called by the host. Here is code example
+of the SYCL program, which demonstrates the need for this attribute:
----------------
is code example -> is a code example


================
Comment at: clang/include/clang/Basic/AttrDocs.td:261
+for the device and cannot be directly called by the host. Here is code example
+of the SYCL program, which demonstrates the need for this attribute:
+
----------------
the SYCL program, which -> a SYLC program that

(Note, I also removed the comma.)


================
Comment at: clang/include/clang/Basic/AttrDocs.td:277
+
+Code is passed to parallel_for is called "kernel function" and defines some
+entry point to device code i.e. will be called by host in run time. Compiler
----------------
Code -> Do you mean the lambda? If so, perhaps "The lambda that is passed to the `parallel_for`" (add backticks around parallel_for too, please).

is called "kernel function" -> is called a "kernel function"


================
Comment at: clang/include/clang/Basic/AttrDocs.td:286
+help.
+  }];
+}
----------------
I'm still not entirely certain how I would know what to mark and how. From the description, it sounds like whoever authors `parallel_for` needs to do this marking, or it somehow happens automatically?

(I'll do another editorial pass once I understand the intended behavior a bit better -- I expect there will be a few more wording issues to address.)


================
Comment at: clang/test/SemaSYCL/device-attributes-on-non-sycl.cpp:2
+// RUN: %clang_cc1 -fsyntax-only -fsycl-is-device -verify %s
+// Now pretend that we're compiling a C++ file. There should be warnings.
+// RUN: %clang_cc1 -fsyntax-only -verify -x c++ %s
----------------
This comment confuses me -- the file *is* a C++ file and the preceding RUN line treats it as such, doesn't it?


================
Comment at: clang/test/SemaSYCL/device-attrubutes.cpp:8-11
+__attribute((sycl_kernel)) void foo();
+__attribute((sycl_device)) void foo1();
+[[clang::sycl_kernel]] void foo2();
+[[clang::sycl_device]] void foo3();
----------------
bader wrote:
> Fznamznon wrote:
> > bader wrote:
> > > This duplicates clang/test/SemaSYCL/device-attributes-on-non-sycl.cpp.
> > > Maybe it make sense to test case when both attributes are applied to the same function.
> > From current documentation: sycl_kernel can be applied to function which can be directly called by the host and will be compiled for device, sycl_device applies to device functions which **cannot** be directly called by the host... 
> > I think if we want add this test case we need clarify how this case will be handled by the compiler.
> > From current documentation: sycl_kernel can be applied to function which can be directly called by the host and will be compiled for device, sycl_device applies to device functions which cannot be directly called by the host... 
> > I think if we want add this test case we need clarify how this case will be handled by the compiler.
> 
> I would expect `sycl_kernel` to supersede `sycl_device` attribute. Basically `sycl_kernel` include functionality of `sycl_device` attribute and provides additional host accessibility.
> I would expect sycl_kernel to supersede sycl_device attribute. Basically sycl_kernel include functionality of sycl_device attribute and provides additional host accessibility.

I'm not a SYCL person, but I would not expect that at all based on the description above. I would expect it to be an error -- it sounds like these attributes are intended to be mutually exclusive?


Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D60455/new/

https://reviews.llvm.org/D60455