keryell 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}}
----------------
bader wrote:
> 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?
>
> If I understand you concerns correctly, you want to be sure that clang
> prohibits other uses of this attribute, which are not intended. Right?
But since it is an attribute to be used by SYCL run-time writers, I am not sure
there is a lot of value in over-engineering the restrictions of its use. It
diverts brain power from the real implementation & review and might even
prevent innovation due to some creative use cases.
Repository:
rG LLVM Github Monorepo
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D60455/new/
https://reviews.llvm.org/D60455
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
