[PATCH] D60455: [SYCL] Implement SYCL device code outlining
Aaron Ballman via Phabricator via cfe-commits
cfe-commits at lists.llvm.org
Mon Jun 24 13:28:39 PDT 2019
aaron.ballman 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}}
----------------
keryell wrote:
> 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.
>
>
> If I understand you concerns correctly, you want to be sure that clang prohibits other uses of this attribute, which are not intended. Right?
Effectively, yes. I'd like to ensure that situations where the attribute does not do what the user expects are diagnosed. A good rule of thumb that I use is to diagnose (as a warning) situations where the attribute will be silently ignored, and diagnose (as an error) situations where applying the attribute would cause really bad results (like miscompiles, security concerns, etc).
> What is the best way to do this? Add more negative tests cases and make sure that clang generate error diagnostic messages?
That's a good approach, yes. Though for the situations you describe, I'd probably just warn rather than err because it seems like it's harmless to ignore the attribute so long as the user knows it's being ignored.
================
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:
> keryell wrote:
> > 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.
> >
> >
> > If I understand you concerns correctly, you want to be sure that clang prohibits other uses of this attribute, which are not intended. Right?
>
> Effectively, yes. I'd like to ensure that situations where the attribute does not do what the user expects are diagnosed. A good rule of thumb that I use is to diagnose (as a warning) situations where the attribute will be silently ignored, and diagnose (as an error) situations where applying the attribute would cause really bad results (like miscompiles, security concerns, etc).
>
> > What is the best way to do this? Add more negative tests cases and make sure that clang generate error diagnostic messages?
>
> That's a good approach, yes. Though for the situations you describe, I'd probably just warn rather than err because it seems like it's harmless to ignore the attribute so long as the user knows it's being ignored.
> 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.
I disagree. Part of the real implementation is ensuring the attribute is not accidentally misused. It's frustrating for users to have an attribute silently ignored because it's easy to mistake that situation for the attribute behaving as expected.
Repository:
rG LLVM Github Monorepo
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D60455/new/
https://reviews.llvm.org/D60455
More information about the cfe-commits
mailing list