[PATCH] D60455: [SYCL] Implement SYCL device code outlining

[Alexey Bader via Phabricator via cfe-commits]

bader added inline comments.


================
Comment at: clang/test/SemaSYCL/device-attributes.cpp:3
+
+[[clang::sycl_kernel]] int gv2 = 0; // expected-warning {{'sycl_kernel' attribute only applies to functions}}
+__attribute((sycl_kernel)) int gv3 = 0; // expected-warning {{'sycl_kernel' attribute only applies to functions}}
----------------
aaron.ballman wrote:
> Fznamznon wrote:
> > aaron.ballman wrote:
> > > I'd like to see some more tests covering less obvious scenarios. Can I add this attribute to a lambda? What about a member function? How does it work with virtual functions? That sort of thing.
> > Actually there is no restrictions for adding this attribute to any function to outline device code so I just checked the simplest variant.
> > 
> > But I'm working on new patch which will put some requirements on function which is marked with `sycl_kernel` attribute. 
> > This new patch will add generation of OpenCL kernel from function marked with `sycl_kernel` attribute. The main idea of this approach is described in this [[ https://github.com/intel/llvm/blob/sycl/sycl/doc/SYCL_compiler_and_runtime_design.md#lowering-of-lambda-function-objects-and-named-function-objects | document ]] (in this document generated kernel is called "kernel wrapper").
> > And to be able to generate OpenCL kernel using function marked with `sycl_kernel` attribute we put some requirements on this function, for example it must be a template function. You can find these requirements and example of proper function which can be marked with `sycl_kernel` in this [[ https://github.com/intel/llvm/pull/177#discussion_r290451286 | comment ]] .
> > 
> > 
> > Actually there is no restrictions for adding this attribute to any function to outline device code so I just checked the simplest variant.
> 
> So there are no concerns about code like:
> ```
> struct Base {
>   __attribute__((sycl_kernel)) virtual void foo();
>   virtual void bar();
> };
> 
> struct Derived : Base {
>   void foo() override;
>   __attribute__((sycl_kernel)) void bar() override;
> };
> 
> void f(Base *B, Derived *D) {
>   // Will all of these "do the right thing"?
>   B->foo();
>   B->bar();
> 
>   D->foo();
>   D->bar();
> }
> ```
> Actually there is no restrictions for adding this attribute to any function to outline device code so I just checked the simplest variant.
> But I'm working on new patch which will put some requirements on function which is marked with sycl_kernel attribute.

@aaron.ballman, sorry for confusing. The  usage scenarios should have been articulated more accurately.
We have only four uses of this attribute in our implementation:
https://github.com/intel/llvm/blob/sycl/sycl/include/CL/sycl/handler.hpp#L538 (lines 538-605).
All four uses are applied to member functions of `cl::sycl::handler` class and all of them have similar prototype (which is mentioned by Mariya in the previous [comment](https://github.com/intel/llvm/pull/177#discussion_r290451286):

```
namespace cl { namespace sycl {
class handler {
  template <typename KernelName, typename KernelType/*, ...*/>
  __attribute__((sycl_kernel)) void sycl_kernel_function(KernelType KernelFuncObj) {
    KernelFuncObj();
  }
};
}}
```

Here is the list of SYCL device compiler expectations with regard to the function marked with `sycl_kernel` attribute.
	- Function template with at least one parameter is expected. The compiler generates OpenCL kernel and uses first template parameter as unique name to the generated OpenCL kernel. Host application uses this unique name to invoke the OpenCL kernel generated for the `sycl_kernel_function` specialized by this name and KernelType (which might be a lambda type).
	- Function must have at least one parameter. First parameter expected to be a function object type (named or unnamed i.e. lambda). Compiler uses function object type field to generate OpenCL kernel parameters.

Aaron, I hope it makes more sense now.

We don't plan in any use cases other than in SYCL standard library implementation mentioned above.
If I understand you concerns correctly, you want to be sure that clang prohibits other uses of this attribute, which are not intended. Right?
What is the best way to do this? Add more negative tests cases and make sure that clang generate error diagnostic messages?



Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D60455