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
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
