[SYCL] Decompose kernel parameters and add inheritance support #1877

elizabethandrews · 2020-06-12T04:47:51Z

This PR redesigns 'parameter passing' support for kernel parameters of struct type. All struct type SYCL kernel arguments (except special SYCL types) are now decomposed and their individual fields are passed as separate OpenCL kernel arguments.

Following issues are fixed after this PR:

Inheritance support for SYCL Kernel
Accessors in base class
Performance issues due to passing SYCL special types twice.
Nested Arrays

This PR was a collaborative effort, with patches from Mariya Podchishchaeva and Elizabeth Andrews. For early testing, PR was updated periodically with changes from PR#1841 (Array support by Rajiv Deodhar). Changes from PR#1841 is not present in final diff since array support is now in intel:sycl.

Signed-off-by: rdeodhar <rajiv.deodhar@intel.com>

Signed-off-by: Mariya Podchishchaeva <mariya.podchishchaeva@intel.com>

Vector types started appearing as top level arguments since they are wrapped with cl::sycl::vec class. Signed-off-by: Mariya Podchishchaeva <mariya.podchishchaeva@intel.com>

…They are decomposed and individual fields are passed Signed-off-by: Elizabeth Andrews <elizabeth.andrews@intel.com>

Signed-off-by: Elizabeth Andrews <elizabeth.andrews@intel.com>

through it's fields and initialize accessors using respective handlers. Other fields of stream class are not handled.

elizabethandrews · 2020-06-12T05:03:27Z

I think all lit tests except for those containing arrays should now pass. I think there are 3 with arrays in them.

elizabethandrews · 2020-06-12T16:43:20Z

There is an issue with multiple streams in one command group. Debugging that now.

Signed-off-by: Elizabeth Andrews <elizabeth.andrews@intel.com>

iburyl · 2020-06-15T09:23:37Z

When I need to allow inherited accessors in my own branch, I need to change those two functions (clang/lib/Sema/SemaSYCL.cpp):

getAccessTarget - this one expected target template argument to be exactly the fourth one
isSyclAccessorType - this one expected the class of accessor to be named "accessor"

Are you sure, you do need changing those?

Fznamznon · 2020-06-15T11:30:28Z

When I need to allow inherited accessors in my own branch, I need to change those two functions (clang/lib/Sema/SemaSYCL.cpp):

getAccessTarget - this one expected target template argument to be exactly the fourth one

isSyclAccessorType - this one expected the class of accessor to be named "accessor"

This approach won't work for user-defined classes which have accessor as base class. We need to handle this generic case as well.

Are you sure, you do need changing those?

We don't changing those.
If you mean that we do need change those, I'm not sure about it because of the reasons I explained above.

Fznamznon · 2020-06-15T11:36:14Z

We don't changing those.
If you mean that we do need change those, I'm not sure about it because of the reasons I explained above.

I think to be sure here, we need to understand what is the purpose of your case. If your inherited accessor has its own implementation of __init method and you want it to be a top level kernel argument like regular accessor now, you may need to change those functions.

Fznamznon · 2020-06-15T11:41:53Z

clang/lib/Sema/SemaSYCL.cpp

+// FIXME: Can this be refactored/handled some other way?
+template <typename ParentTy, typename... Handlers>
+static void VisitStreamRecord(CXXRecordDecl *Owner, ParentTy &Parent,
+                              CXXRecordDecl *Wrapper, Handlers &... handlers) {
  (void)std::initializer_list<int>{(handlers.enterStruct(Owner, Parent), 0)...};
-  VisitAccessorWrapperHelper(Wrapper, Wrapper->bases(), handlers...);
-  VisitAccessorWrapperHelper(Wrapper, Wrapper->fields(), handlers...);
+  for (const auto &Field : Wrapper->fields()) {
+    QualType FieldTy = Field->getType();
+    // Required to initialize accessors inside streams.
+    if (Util::isSyclAccessorType(FieldTy))
+      KF_FOR_EACH(handleSyclAccessorType, Field, FieldTy);
+  }
  (void)std::initializer_list<int>{(handlers.leaveStruct(Owner, Parent), 0)...};
 }


I think the best way for us is to re-implement stream class in SYCL headers, so it will look like accessor and we don't need to visit it's field in the compiler.
@againull , do you think it is possible?

clang/lib/Sema/SemaSYCL.cpp

iburyl · 2020-06-15T14:45:21Z

We don't changing those.
If you mean that we do need change those, I'm not sure about it because of the reasons I explained above.

I think to be sure here, we need to understand what is the purpose of your case. If your inherited accessor has its own implementation of __init method and you want it to be a top level kernel argument like regular accessor now, you may need to change those functions.

Our use case is having some_accessor, which is inherited from accessor.
We have no problem writing own __init (though avoiding that would be ideal).
Now this some_accessor is passed to kernel.

Will this use case be covered by your patch?

My own approach was: writing __init, changing getAccessTarget, for it to parse type parents, until it finds accessor, then check target template argument there + isSyclAccessorType is changed, to make some_accessor known to it. That worked.

@Pennycook

Fznamznon · 2020-06-15T15:37:57Z

Our use case is having some_accessor, which is inherited from accessor.
We have no problem writing own __init (though avoiding that would be ideal).
Now this some_accessor is passed to kernel.

Our intention was to implement support for user-defined inherited accessors. So passing some class which have accessor in base should work and __init writing is not required.

…upport_2

Signed-off-by: Elizabeth Andrews <elizabeth.andrews@intel.com>

erichkeane

NO comments on this for me, the stuff I understand all looks right.

rdeodhar

All the changes look good to me.

…upport_2 Signed-off-by: Elizabeth Andrews <elizabeth.andrews@intel.com>

Signed-off-by: Elizabeth Andrews <elizabeth.andrews@intel.com>

elizabethandrews · 2020-07-02T23:19:03Z

Merged latest changes from intel:sycl, fixing conflict (Just a variable name change).

Fznamznon · 2020-07-03T09:21:33Z

@intel/llvm-reviewers-runtime there are some runtime tests, could someone review them, please?

s-kanaev

The runtime tests look good to me, except for a single comment there.
Also, I believe it's an extension and not by-specification. Hence, a link to the extension proposal/spec would be a nice to have.

sycl/test/array_param/array-kernel-param-nested-run.cpp

Fznamznon · 2020-07-03T14:24:10Z

Also, I believe it's an extension and not by-specification. Hence, a link to the extension proposal/spec would be a nice to have.

This is not an extension. SYCL specification technically doesn't prohibit inheritance for kernel arguments, kernel function objects and accessors, so this patch only adds handling of these missed cases.

s-kanaev · 2020-07-03T14:26:24Z

This is not an extension. SYCL specification technically doesn't prohibit inheritance for kernel arguments, kernel function objects and accessors, so this patch only adds handling of these missed cases.

OK, then.

bader

A couple of improvement suggestions, which can be applied in separate PRs.

bader · 2020-07-06T12:57:39Z

clang/test/CodeGenSYCL/struct_kernel_param.cpp

+// CHECK-NEXT:  { kernel_param_kind_t::kind_std_layout, 4, 12 },
+// CHECK-NEXT:  { kernel_param_kind_t::kind_std_layout, 4, 16 },
+// CHECK-NEXT:  { kernel_param_kind_t::kind_std_layout, 4, 20 },
+// CHECK-NEXT:  { kernel_param_kind_t::kind_std_layout, 4, 24 },
+// CHECK-NEXT:  { kernel_param_kind_t::kind_std_layout, 4, 28 },
+// CHECK-NEXT:  { kernel_param_kind_t::kind_std_layout, 4, 32 },


I don't think we need to decompose a struct consisting of standard trivially-copiable C++ types (i.e. which doesn't encapsulate a non-USM pointer or non-standard C++ type e.g. image or sampler).
Sending it like this will impact performance:

we set 6 parameters instead of one, which means runtime overhead on calling "setArg" 6 times + potentially use more registers than needed to pass six 4-byte values instead of one 24-byte value

it might complicate JIT compiler analysis (or may not).

I'm okay if we address this in follow-up commit.

Yes, we were planning to add this optimization after implementing basic functionality. I've created an issue to keep track on that #2043. @elizabethandrews , would you mind taking this one?

Sure will do.

bader · 2020-07-06T13:18:04Z

clang/lib/Sema/SemaSYCL.cpp

+      Util::DeclContextDesc{clang::Decl::Kind::Namespace, "cl"},
+      Util::DeclContextDesc{clang::Decl::Kind::Namespace, "sycl"},
+      Util::DeclContextDesc{clang::Decl::Kind::Namespace, "detail"},
+      Util::DeclContextDesc{clang::Decl::Kind::Namespace, "half_impl"},
+      Util::DeclContextDesc{Decl::Kind::CXXRecord, Name}};


Please, add an attribute to annotate SYCL types like accessor, sampler, image. It's much more reliable approach than declaration name, which introduces hard-coded dependency on exact symbol name in the runtime library.
I suppose we need only one attribute to detect special types which must be handled as separate kernel parameters. If not, we can parameterize the attribute.

This refactoring can be done in a separate PR.
Tagging @Naghasan, who was going to contribute similar attribute from ComputeCPP implementation.

Please, add an attribute to annotate SYCL types like accessor, sampler, image. It's much more reliable approach than declaration name, which introduces hard-coded dependency on exact symbol name in the runtime library.
I suppose we need only one attribute to detect special types which must be handled as separate kernel parameters. If not, we can parameterize the attribute.

This refactoring can be done in a separate PR.
Tagging @Naghasan, who was going to contribute similar attribute from ComputeCPP implementation.

Right now just one attribute is not enough since we have differences between handling of different classes, especially in integration header support. It seems runtime library needs to handle each type in a different way, I have no idea how to redo it so it will be unified across different SYCL types.
Right now we can start with parameterized attribute, i.e. for example:

class accessor { ... __attribute__((sycl_special_type(accessor))) DataT *Data; } ... class sampler { ... __attribute__((sycl_special_type(sampler))) ocl_sampler_t Sampler; }

Also, we can discuss it in upcoming upstreaming meeting.
I've created an issue to keep track on that #2041.

I'd like to point out that clang is able to differentiate sampler from accessor w/o any attributes - they are different types (pointer type vs __ocl_sampler_t).

Okay, makes sense.

This type instruction describes a string, mostly for Fortran 90. Spec: KhronosGroup/SPIRV-Registry#186 Original commit: KhronosGroup/SPIRV-LLVM-Translator@beecd9d

rdeodhar and others added 11 commits June 8, 2020 17:04

Support for arrays as kernel parameters.

7974c35

Signed-off-by: rdeodhar <rajiv.deodhar@intel.com>

Reusing some memberexpr building code.

4907194

Merge branch 'sycl' of https://github.com/otcshare/llvm into akp2

d54c0ca

Rebased changes from Elizabeth

44d8663

Signed-off-by: Mariya Podchishchaeva <mariya.podchishchaeva@intel.com>

Implement special bases handling

801a0ea

Signed-off-by: Mariya Podchishchaeva <mariya.podchishchaeva@intel.com>

Merge branch 'sycl' into bases-handling

503638e

Merge branch 'sycl' of https://github.com/otcshare/llvm into akp2

546c58d

Handle vector types like scalar types

52f2e5a

Vector types started appearing as top level arguments since they are wrapped with cl::sycl::vec class. Signed-off-by: Mariya Podchishchaeva <mariya.podchishchaeva@intel.com>

Fixed failing lit tests. Structs/Classes are no longer passed whole. …

2a36a93

…They are decomposed and individual fields are passed Signed-off-by: Elizabeth Andrews <elizabeth.andrews@intel.com>

Owner should be record being visited.

ab74fcf

Signed-off-by: Elizabeth Andrews <elizabeth.andrews@intel.com>

Avoid decomposing stream class. If field type is stream, we iterate

4370d76

through it's fields and initialize accessors using respective handlers. Other fields of stream class are not handled.

elizabethandrews requested review from Fznamznon and premanandrao as code owners June 12, 2020 04:47

rdeodhar and others added 5 commits June 12, 2020 10:41

Updated support for arrays.

52ce3f2

Merge branch 'sycl' of https://github.com/otcshare/llvm into akp2

983b3d5

Formatting changes.

1bf0903

Formatting changes.

5d5121b

Fix crash for stream type. Also changed handler call for consistency

96ca8f4

Signed-off-by: Elizabeth Andrews <elizabeth.andrews@intel.com>

Fznamznon reviewed Jun 15, 2020

View reviewed changes

rdeodhar and others added 4 commits June 15, 2020 09:00

Correction to a test.

f03edd9

Merge remote-tracking branch 'rajiv_fork/akp2' into add_inheritance_s…

033b507

…upport_2

Fixed some crashes after merge.

4868d45

Signed-off-by: Elizabeth Andrews <elizabeth.andrews@intel.com>

Minor refactor

35383c5

Signed-off-by: Elizabeth Andrews <elizabeth.andrews@intel.com>

elizabethandrews changed the title ~~[WIP][SYCL][Do not review]Decompose kernel parameters and inheritance support~~ Decompose kernel parameters and add inheritance support Jul 2, 2020

elizabethandrews marked this pull request as ready for review July 2, 2020 18:03

elizabethandrews requested review from erichkeane and rdeodhar July 2, 2020 18:04

elizabethandrews changed the title ~~Decompose kernel parameters and add inheritance support~~ [SYCL] Decompose kernel parameters and add inheritance support Jul 2, 2020

erichkeane reviewed Jul 2, 2020

View reviewed changes

erichkeane previously approved these changes Jul 2, 2020

View reviewed changes

rdeodhar previously approved these changes Jul 2, 2020

View reviewed changes

premanandrao previously approved these changes Jul 2, 2020

View reviewed changes

elizabethandrews added 2 commits July 2, 2020 16:12

Merge remote-tracking branch 'intel_llvm/sycl' into add_inheritance_s…

47d092a

…upport_2 Signed-off-by: Elizabeth Andrews <elizabeth.andrews@intel.com>

ClangFormat Change

a7ad39c

Signed-off-by: Elizabeth Andrews <elizabeth.andrews@intel.com>

elizabethandrews dismissed stale reviews from premanandrao, rdeodhar, and erichkeane via a7ad39c July 2, 2020 23:16

elizabethandrews requested a review from premanandrao July 2, 2020 23:19

This was linked to issues Jul 3, 2020

[SYCL] Cannot use derived accessor objects #28

Closed

[SYCL][SemaSYCL] Functor Inheritance Incorrect Outline/CodeGen #488

Closed

s-kanaev approved these changes Jul 3, 2020

View reviewed changes

sycl/test/array_param/array-kernel-param-nested-run.cpp Show resolved Hide resolved

Fznamznon approved these changes Jul 6, 2020

View reviewed changes

bader linked an issue Jul 6, 2020 that may be closed by this pull request

[SYCL] Confirm support for nested arrays after struct decomposition complete in PR#1877 #2018

Closed

bader reviewed Jul 6, 2020

View reviewed changes

bader merged commit 0b2de9e into intel:sycl Jul 6, 2020

This was referenced Jul 6, 2020

[SYCL][FE] Add an attribute to annotate special sycl types #2041

Closed

[SYCL][FE] Decompose structs on demand #2043

Closed

bader mentioned this pull request Jul 10, 2020

[WIP][SYCL] Implement sycl_special_class attribute #2091

Closed

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

[SYCL] Decompose kernel parameters and add inheritance support #1877

[SYCL] Decompose kernel parameters and add inheritance support #1877

elizabethandrews commented Jun 12, 2020 •

edited

Loading

elizabethandrews commented Jun 12, 2020

elizabethandrews commented Jun 12, 2020

iburyl commented Jun 15, 2020

Fznamznon commented Jun 15, 2020

Fznamznon commented Jun 15, 2020

Fznamznon Jun 15, 2020

iburyl commented Jun 15, 2020

Fznamznon commented Jun 15, 2020

erichkeane left a comment

rdeodhar left a comment

elizabethandrews commented Jul 2, 2020

Fznamznon commented Jul 3, 2020

s-kanaev left a comment

Fznamznon commented Jul 3, 2020 •

edited

Loading

s-kanaev commented Jul 3, 2020

bader left a comment

bader Jul 6, 2020

Fznamznon Jul 6, 2020

elizabethandrews Jul 6, 2020

bader Jul 6, 2020

Fznamznon Jul 6, 2020

bader Jul 6, 2020

Fznamznon Jul 6, 2020

[SYCL] Decompose kernel parameters and add inheritance support #1877

[SYCL] Decompose kernel parameters and add inheritance support #1877

Conversation

elizabethandrews commented Jun 12, 2020 • edited Loading

elizabethandrews commented Jun 12, 2020

elizabethandrews commented Jun 12, 2020

iburyl commented Jun 15, 2020

Fznamznon commented Jun 15, 2020

Fznamznon commented Jun 15, 2020

Choose a reason for hiding this comment

iburyl commented Jun 15, 2020

Fznamznon commented Jun 15, 2020

erichkeane left a comment

Choose a reason for hiding this comment

rdeodhar left a comment

Choose a reason for hiding this comment

elizabethandrews commented Jul 2, 2020

Fznamznon commented Jul 3, 2020

s-kanaev left a comment

Choose a reason for hiding this comment

Fznamznon commented Jul 3, 2020 • edited Loading

s-kanaev commented Jul 3, 2020

bader left a comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

elizabethandrews commented Jun 12, 2020 •

edited

Loading

Fznamznon commented Jul 3, 2020 •

edited

Loading