[WIP][SYCL] Implement sycl_special_class attribute

[Fznamznon]

This attribute is used in SYCL headers to mark SYCL classes which need
additional compiler handling when passed from host to device.
Attribute can be applied to struct/class and can have optional argument
which indicates kind of SYCL special class, so it can be used to
implement handling of some generic case of SYCL special class as well as
implement different handling for each kind of SYCL special class.

Usage:

class __attribute__((sycl_special_class(accessor))) accessor {
...
}

The PR is created to gather early community feedback regarding the attribute form, since similar approach was discussed on upstreaming WG.

To finish this one, the following needs to be done:
-[ ] Add documentation
-[ ] Add a test for attribute
-[ ] Resolve TODO about spec constant class detection

Closes #2041

[Fznamznon]

Tagging @Naghasan, who was going to contribute similar attribute from ComputeCPP implementation.

[bader]

The PR is created to gather early community feedback regarding the attribute form, since similar approach was discussed on upstreaming WG.

That implementation is different from what was discussed here: #1877 (comment).

I think this might be OpenCL/SPIR-V specific attribute instead of SYCL specific attribute as it's supposed to address OpenCL requirements for kernel parameters that non-USM pointers, samplers, images, pipes must be separate kernel parameters and can't be used in aggregate types.

I think it would make more sense to apply this attribute to the objects, which we should pass as a separate kernel parameter.

class <some class> {
  ...
  [[opencl::kernel_parameter]] <type> <member name>;
  <another type> <another member name>;
}

SYCL compiler can apply the same logic to all object of types where this attribute is applied. Pass members annotated with the attribute as a separate kernel parameter, pass other member of the class with annotated members separately.

class accessor {
  ...
  [[opencl::kernel_parameter]] DataT *Data;
}

class sampler {
  ...
  [[opencl::kernel_parameter]] ocl_sampler_t Sampler;
}

[Naghasan]

I have mainly looked at the attribute definition here.

This seems to me to be very specific to SYCL and does not cover every special types (like specialization constant). Also if you want to add a new one, you need to change the compiler.

One of the idea I had was to make the attribute sensitive to __init and __finalize (if present) and maintain the derivation of the kernel arguments from the __init argument list. This allow to add new special types without the need to touch the compiler and you are actually SYCL agnostic. As this is a special class, the runtime should be able to query what the class is and act accordingly. I never had the time to fully explore and test the idea, so it might be flawed. There is also the case of the specialization constant, where you may or may not want to pass the value as an argument, but maybe this could simply be handled using macros.

One drawback I see from a "generic attribute" point of view, this is C++ specific as __init and __finalize are expected to be member functions. But the functions could be specify as part of the argument.

I think this might be OpenCL/SPIR-V specific attribute instead of SYCL specific attribute

I disagree, OpenCL doesn't need to apply any kind of processing to build the entry point. SPIR-V is neutral to that regard as well.

This is to allow regularization to an underlying programming model. So OpenCL or CUDA at the moment, but we could very well imagine something OpenMP or Vulkan compatible.

[Fznamznon]

This seems to me to be very specific to SYCL and does not cover every special types (like specialization constant). Also if you want to add a new one, you need to change the compiler.

Yes, it doesn't cover specialization constant, I marked it as unfinished TODO in the PR description. In case of new special type... If the new special type needs some special handling that differs from others, you need to change the compiler anyway, otherwise "generic" value of the attribute can be used.

One of the idea I had was to make the attribute sensitive to __init and __finalize (if present) and maintain the derivation of the kernel arguments from the __init argument list.

Could you please describe how it would look like? I'm not sure that I got the idea.

As this is a special class, the runtime should be able to query what the class is and act accordingly.

With the current implementation the runtime takes the address of kernel object and uses offsets encoded by the integration header. It also uses information about kernel object field types to cast area of memory accessed by using of address and offset to proper type and perform handling, it cannot query the type from raw memory. In this case we probably could teach integration header to emit information that some kernel argument is a 'special class' and add some base class for each class that is handled specially by the compiler, so this base class will hold kind of special class and it will be possible to query it, I guess. But I'm not really sure since I don't have a lot of experience with runtime.

[Naghasan]

If the new special type needs some special handling that differs from others, you need to change the compiler anyway, otherwise "generic" value of the attribute can be used.

This attribute doesn't do much. Special types processing boils down to fully or partially breakdown the structure into valid fields for your programming model. The rest is the associated processing in the runtime.

Could you please describe how it would look like? I'm not sure that I got the idea.

I'll take the accessor as an example.
ATM (and last time I checked) the processing of the special type sycl::accessor is roughly done as follow:

• Check if sycl special type sycl::accessor, if yes
• Look up for the member function sycl::accessor::__init
• Copy the sycl::accessor::__init arguments into the kernel function argument list
• Allocate a sycl::accessor object and call sycl::accessor::__init with the kernel argument

This is spread out in different places, but that's the basic idea. The important thing to notice is if you replace sycl::accessor by sycl::sampler, sycl::stream or other, the processing remains the same. The type sycl::stream has an extra step which is to call __finalize at the end of the kernel.

So in a more generic way, you could have __attribute__((sycl_special_class)) (leaving the renaming exercise for later) that actually trigger the process above in a more generic way:

• Check if <typename> has attribute sycl_special_class, if yes
• Look up for the member function <typename>::__init
• Add kernel function has friend of the class
• Copy the <typename>::__init arguments into the kernel function argument list
• Allocate a <typename> object and call <typename>::__init with the kernel argument
• Look up for the member function <typename>::__finalize, if does have the function
• Call <typename>::__finalize at the end of the kernel

So for example:

class __attribute__((sycl_special_class)) MySpecialType {
  int Field1;
  int Field2;
  void __init(int F1) {
    Field1 = F1;
    Field2 = F1;
  }
  void __finalize() {}
public:
  MySpecialType() = default;
  int getF2() const { return Field2; }
};

If used in a kernel argument

MySpecialType T;
cgh.single_task([=]() {
  T.getF2();
});

This would trigger the following kernel entry point in the AST:

void __sycl_kernel(int F1) {
  MySpecialType T;
  T.__init(F1);
  // finish rebuilding the lambda + call
  T.__finalize()
}

With the current implementation the runtime takes the address of kernel object and uses offsets encoded by the integration header. It also uses information about kernel object field types to cast area of memory accessed by using of address and offset to proper type and perform handling, it cannot query the type from raw memory. In this case we probably could teach integration header to emit information that some kernel argument is a 'special class' and add some base class for each class that is handled specially by the compiler, so this base class will hold kind of special class and it will be possible to query it, I guess. But I'm not really sure since I don't have a lot of experience with runtime.

So the integration header already emits what kind of fields it is processing. Changing to this would only require the capability to query it in the runtime. Since the std-layout requirement was lifted, a simple base class and reinterpret cast should do the job.

Another approach is to make the attribute take an ID as argument (the ID is chosen by the runtime) and forward it in the integration header (so in place of the field currently emitted). The runtime then only have to read-it directly. Either way, the compiler remains agnostic to the field it is processing.

Hopes this make the idea clearer.

[Fznamznon]

Yes, it definitely makes the idea clearer. Thanks for so detailed explanation!
I'll give it a try. It will require more changes across front-end and runtime library, so it will take some time. I will update this PR once I have something ready.

[romanovvlad]

Minor. Suggest adding check: #ifdef SYCL_DEVICE_ONLY

[Fznamznon]

I think it would make more sense to apply this attribute to the objects, which we should pass as a separate kernel parameter.

class <some class> {
  ...
  [[opencl::kernel_parameter]] <type> <member name>;
  <another type> <another member name>;
}

I didn't see the purpose in applying attribute to a member of some SYCL class instead of applying it directly to the class because current implementation combines kernel arguments for each accessor/sampler/stream (i.e. things that I and SYCL spec call special classes) using __init method arguments.
Could you please elaborate why attribute that is applied to some member is better than attribute that is applied to a whole record?

SYCL compiler can apply the same logic to all object of types where this attribute is applied. Pass members annotated with the attribute as a separate kernel parameter, pass other member of the class with annotated members separately.

I didn't get what you mean here. If one member is passed as a separate kernel parameter, the others should be passed in the same way, right? Because we don't wan't to pass some members twice.

[bader]

FYI: @erichkeane recently added SYCLRequiresDecomposition attribute, which probably can be re-used instead of SYCLSpecialClass.
15e62c2

[Fznamznon]

SYCLRequiresDecomposition attribute appears on user-defined struct kernel arguments if they should be passed to the kernel field-by-field, then these structs are constructed back on kernel side without init method usage. Classes which we want to identify with SYCLSpecialClass (accessor, sampler, stream) are handled specially with init method. So, technically right now SYCLRequiresDecomposition attribute works in a bit other use case. To re-use this attribute we will need to change some logic.

[erichkeane]

presumably SYCLSpecialClass is a subset of the SYCLRequiresDecomposition. I think we still would need BOTH, as the Special-class needs to understand how to deal with the 'init' function, but it should also be translated to the requires-decomp.

[bader]

@Fznamznon, what is the plan for this PR?
I would prefer this attribute upstreamed before https://reviews.llvm.org/D71016.

[Fznamznon]

@Fznamznon, what is the plan for this PR?
I would prefer this attribute upstreamed before https://reviews.llvm.org/D71016.

I was planning to align handling of SYCL classes (see #2268) then proceed with this attribute. However since I have no bandwidth anymore, I'm not sure about the plan.

[elizabethandrews]

Can we close this? Its replaced by - #3892