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

[Mariya Podchishchaeva via Phabricator via cfe-commits]

Fznamznon added a comment.

I tried to reuse OpenCL kernel attribute with "__kernel" keyword in our current SYCL implementation. PR with this try is here  - https://github.com/intel/llvm/pull/97
Now It looks feasible but with a couple notes:
>From SYCL specification "SYCL is designed to be as close to standard C++ as possible. Standard C++ compiler can compile the SYCL programs and they will run correctly on host CPU." So SYCL doesn't provide non-standard `kernel` keyword which is provided by OpenCL. Due this fact it's not possible to add `kernel` keyword as in OpenCL, it will prevent compilation of following valid SYCL code:

  int foo(int kernel) { return ++kernel; } // If "kernel" will be a keyword like in OpenCL, here will be a error
  …
  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);
        });
      }
  ...

So I added only `__kernel` keyword for SYCL because in C++ identifiers which start with `__` are reserved for compiler internals.
Next note:
In our current implementation actually not quite that function which is marked with `sycl_kernel` (or `__kernel`, whatever) will be real OpenCL kernel in produced module. In SYCL all shared between host and device memory objects (buffers/images, these objects map to OpenCL buffers and images) can be accessed through special `accessor` classes. SYCL also has special mechanism for passing kernel arguments from host to device, if in OpenCL you need to do `clSetKernelArg`, in SYCL all kernel arguments are captures/fields of lambda/functor which is passed to `parallel_for` (See code snippet above, here one kernel argument - accessor `A` ). To map to OpenCL setting kernel arguments mechanism we added generation of some "kernel wrapper" function inside the compiler. "Kernel wrapper" function contains body of SYCL kernel function, receives OpenCL like parameters and additionally does some manipulation to initialize captured lambda fields with this parameters. In some pseudo code "kernel wrapper" looks like this:

  // SYCL kernel is defined in SYCL headers
  __kernel someSYCLKernel(lambda) {
    lambda();
  }
  // Kernel wrapper
  __kernel wrapper(global int* a) {
    lambda; // Actually lambda declaration doesn't have a name in AST
    // Let lambda has one captured field - accessor A. We need to init it with global pointer from arguments:
    lambda.A.__init(a);
    // Body of SYCL kernel from SYCL headers:
    {
      lambda();
    }
  }

And actually kernel wrapper is presented in result module and passed to OpenCL backend.
As I said, kernel wrapper is generated by the compiler inside the Sema and OpenCLKernel attribute manually added to it, no matter which attribute was added to "SYCL kernel" from SYCL headers.
So, while we are generating this wrapper I see only one profit to use OpenCL kernel attribute in SYCL kernels - don't add new attribute to clang (but we need to add `__kernel` keyword to SYCL).
I thought about idea - don't generate kernel wrapper but looks like it will not work with OpenCL since we can't pass OpenCL `cl_mem` arguments inside any structures (including accessors and lambdas) to the kernel.


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D60455