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

Tue Apr 16 12:49:31 PDT 2019

aaron.ballman added a comment.

In D60455#1468714 <https://reviews.llvm.org/D60455#1468714>, @bader wrote:

> In D60455#1468386 <https://reviews.llvm.org/D60455#1468386>, @Fznamznon wrote:
>
> > > Ok, my question is whether you are planning to duplicate the same logic as for OpenCL kernel which doesn't really seem like an ideal design choice. Is this the only difference then we can simply add an extra check for SYCL compilation mode in this template handling case. The overall interaction between OpenCL and SYCL implementation is still a very big unknown to me so it's not very easy to judge about the implementations details...
> >
> > Of course, if nothing prevents us to re-use OpenCL kernel attribute for SYCL I assume it would be good idea. 
> >  But I'm thinking about the situation with https://reviews.llvm.org/D60454 . If we re-use OpenCL kernel attributes - we affected by OpenCL-related changes and OpenCL-related changes shouldn't violate SYCL semantics. Will it be usable for SYCL/OpenCL clang developers? @bader , what do you think about it?
>
>
> I also think it's worth trying. We should be able to cover "SYCL semantics" with LIT test to avoid regressions by OpenCL related changes. E.g. add a test case checking that -fsycl-is-device option disables restriction on applying `__kernel` to template functions.
>  I want to confirm that everyone is okay to enable `__kernel` keyword for SYCL extension and cover SYCL use cases with additional regression tests. IIRC, on yesterday call, @keryell, said that having SYCL specific attributes useful for separation of concerns.

I'm not comfortable with that decision unless the attribute semantics are sufficiently related to justify it. If we're just going to have a lot of `KernelAttr->isSYCL()` vs `KernelAttr->isOpenCL()` accessor calls, it may make more sense to use separate semantic attributes (even if they share spellings), though then I'd be curious how a user combines OpenCL and SYCL attributes.

In D60455#1468779 <https://reviews.llvm.org/D60455#1468779>, @bader wrote:

> @aaron.ballman sorry for confusion.
>  SYCL specification doesn't require user to annotate "device functions" with an attribute - it says following (from section 6.9.1 SYCL functions and methods linkage, https://www.khronos.org/registry/SYCL/specs/sycl-1.2.1.pdf, page 251):
>
> > The default behavior in SYCL applications is that all the definitions and declarations of the functions and methods
> >  are available to the SYCL compiler, in the same translation unit. When this is not the case, all the symbols that
> >  need to be exported to a SYCL library or from a C++ library to a SYCL application need to be defined using the
> >  macro: SYCL_EXTERNAL.
> >  The SYCL_EXTERNAL macro will only be defined if the implementation supports offline linking. The macro is
> >  implementation-defined, but the following restrictions apply:
> > 
> > • SYCL_EXTERNAL can only be used on functions;
> >  • the function cannot use raw pointers as parameter or return types. Explicit pointer classes must be used instead;
> >  • externally defined functions cannot call a cl::sycl::parallel_for_work_item method;
> >  • externally defined functions cannot be called from a cl::sycl::parallel_for_work_group scope.
> > 
> > The SYCL linkage mechanism is optional and implementation defined.
>
> The idea I had is that to define `SYCL_EXTERNAL` macro as `sycl_device` attribute.

That seems sensible, but then we're also missing extra semantic checks in this patch, such as prohibiting raw pointers, etc. Those should get added along with tests.

> BTW, I noticed that `SYCL_EXTERNAL` puts additional requirements `sycl_device` doesn't meet:
> 
>> • SYCL_EXTERNAL can only be used on functions;
> 
> I think our implementation doesn't have such limitations and able to support more use cases.

Does it make sense to deviate from the SYCL spec though?

> Anyway, we can make `sycl_device` attribute implicit for now and return to the implementation of cross translation unit dependencies later.

That might be an easier first-pass for the attribute. I'm not at all opposed to giving it a name, I was just trying to keep the threads of discussion straight. :-)

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

> In D60455#1467279 <https://reviews.llvm.org/D60455#1467279>, @aaron.ballman wrote:
>
> >
>
>
>
>
> > 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?
>
> Yes, we are going to teach the compiler differ SYCL device code from host code and ignore functions without device attributes when SYCL device mode is enabled. I've described some possible implementation details in this comment <https://reviews.llvm.org/D60455#1468386>.

Ah, thank you for that. My preference is to not accept patches introducing attributes without also including the semantic effects from the attributes unless there are extenuating circumstances. Adding the semantics at the same time helps to ensure a more solid code review because it often spawns other design discussion. Do you plan to introduce the semantics in a later version of this patch?

================
Comment at: clang/include/clang/Basic/AttrDocs.td:286
+help.
+  }];
+}
----------------
Fznamznon wrote:
> Anastasia wrote:
> > keryell wrote:
> > > aaron.ballman wrote:
> > > > 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.)
> > > In normal SYCL it happens automatically.
> > > In the compiler unit-tests it is done manually to exercise the framework.
> > > I am the one who suggested that in some other contexts, it could be used manually for some special purpose like using some weird hardware, but I do not want to derail the main review with this.
> > > In normal SYCL it happens automatically.
> > > In the compiler unit-tests it is done manually to exercise the framework.
> > > 
> > 
> > 
> > I think if they are not to be exposed to the user they should have no spelling. There are plenty of other ways to test this. For example AST dump.
> > 
> > 
> > > I am the one who suggested that in some other contexts, it could be used manually for some special purpose like using some weird hardware, but I do not want to derail the main review with this.
> > 
> > If you are suggesting to expose this feature then you are starting some sort of a language extensions and its use should be documented in some way. I am not sure about this but I think we will end up with some sort of a language extension for SYCL anyways because as it stands now it's not aligned with the general concept of C/C++ language design. So perhaps it's not entirely unreasonable to expose this.
> Generally the `sycl_device` attribute will be added automatically by the compiler. But as @bader mentioned before:
> > we might need to use sycl_device attribute to mark functions, which are called from the different translation units, i.e. compiler can't identify it w/o user's help.
> > SYCL specification proposes to use special macro as "device function marker", but I guess we can have additional "spellings" in the clang.
> I think It would be better to re-use this attribute in implementation of "device function marker" macro from SYCL spec than implement additional logic in the compiler to handle this macro. So I saved possibility to add this attribute in code.
> 
> 
> 
> 
I'm not opposed to adding the attribute, but @bader also said that SYCL is not supposed to expose non-standard extensions to users and "we might need this" didn't seem like a strong case for needing the attribute.

If the SYCL spec has a macro that is used to mark the user's code, then 1) that seems like a nonstandard extension to the language, but 2) it does make a strong case for having a named attribute.

================
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
----------------
Fznamznon wrote:
> aaron.ballman wrote:
> > This comment confuses me -- the file *is* a C++ file and the preceding RUN line treats it as such, doesn't it?
> I've tried to say that we compile regular C++ file without SYCL mode enabled. Looks like it can confuse since files with SYCL code have cpp extension... I will rewrite this comment.
Ah, that makes more sense, thanks!

Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D60455