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

[Alexey Bader via Phabricator via cfe-commits]

bader added a comment.

To give more context to the question, I'd like to clarify the use case of new attributes.

As @Fznamznon already mentioned in previous comments SYCL is not supposed to expose any non-standard extensions to a user.

Here is code example of the SYCL program, which demonstrate the need for these attributes:

  C++
  int foo(int x) { return ++x; }
  int bar(int x) { throw std::exception("CPU code only!"); }
  …
  using namespace cl::sycl;
  queue Q;
  buffer<int, 1> a(range<1>{1024});
  Q.submit([&](handler& cgh) {
        auto A = a.get_access<access::mode::write>(cgh);
        cgh.parallel_for<init_a>(range<1>{1024}, [=](id<1> index) {
          A[index] = index[0] * 2 + index[1] + foo(42);
        });
      }
  ...

SYCL compiler need to compile lambda function passed to `cl::sycl::handler::parallel_for` method and function `foo` called from this lambda function.

NOTE: compiler must ignore `bar` function when we "device" part of the single source code.

Our current approach is to add an attribute, which SYCL runtime will use to mark code passed to `cl::sycl::handler::parallel_for` as "kernel functions". Obviously runtime library can't mark `foo` as "device" code - this is a compiler job: to traverse all symbols accessible from kernel functions and add them to the "device part" of the code.

Here is a link to the code in the SYCL runtime using `sycl_kernel` attribute: https://github.com/intel/llvm/blob/sycl/sycl/include/CL/sycl/handler.hpp#L267

I'm quite sure something similar should happen for other "single source" programming models like OpenMP/CUDA, except these attributes are exposed to the user and there is a specific requirement on attributes/pragma/keyword names.

What we are looking for is whether we should add SYCL specific code or we can come up with something more generic to avoid logic/code duplication with already existing functionality.

BTW: Mariya, I think 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.

NOTE2: @Anastasia, https://reviews.llvm.org/D60454 makes impossible to replace `sycl_kernel` attribute with `__kernel` attribute. I mean we still can enable it for SYCL extension, but we will need SYCL specific customization in this case as we apply `kernel` attribute to a template function.


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D60455