[SYCL][ESIMD] Remove wrapping buffer objects into images. #2746

kbobrovs · 2020-11-09T15:09:23Z

This PR removes buffer object wrapping into images, which was used as a SIMD BE workaround and caused a number of problems (e.g. scatter/gather of 1- and 2- byte values worked incorrectly in some cases). It also paves a way for ESIMD/SYCL kernel co-existence in the same source. There are 3 main parts

Runtime

Remove the workaround wrapping.
Gets rid of majority of SYCL_EXPLICIT_SIMD macro checks in the API code
Use the new integration header feature - KernelInfo<>::isESIMD() to properly set accessor arguments
in ESIMD kernels.
Add __init_esimd accessor methods added for FE to pick them up instead of __init for ESIMD kernels, as init sequence is different.

FE

Add 'esimd_acc_ptr' parameter attribute to FE. This attribute is to be used by the FE to mark ESIMD kernel arguments
originating from buffer accessors, which is then translated to BE-specific metadata needed for correct code generation.
Use 'esimd_acc_ptr' attribute to mark ESIMD kernel pointer arguments originating from accessors
Use '__init_esimd' initializer in FE for ESIMD kernel accessors.
Add 'isESIMD()' function to the integration header for the RT to be able to distinguish ESIMD kernels from normal ones for the
purpose of proper accessor arg setting

LowerESIMD

Add 'buffer_t' metadata for accessor kernel arguments and mark it as surface for proper code gen by ESIMD BE.

dm-vodopyanov

ABI is broken. Not sure if we can merge any patches with ABI changes now based on https://github.com/intel/llvm/blob/sycl/sycl/doc/ABIPolicyGuide.md#changing-abi.
@kbobrovs, can this patch be re-written without changes in ABI?

sycl/include/CL/sycl/handler.hpp

DenisBakhvalov · 2020-11-09T23:52:29Z

sycl/include/CL/sycl/accessor.hpp

-  // a global accessor is always wrapped into a 1d image buffer to enable
-  // surface index-based addressing.
-  void __init(OCLImage1dBufferTy ImgBuf) { ImageBuffer = ImgBuf; }
+  const ConcreteASPtrType getNativeImageObj() const { return MData; }


I think having duplicated API is redundant. We can use getQualifiedPtr instead of getNativeImageObj.

That is used in ESIMD APIs. Let's rename/refactor as a separate PR

I added a TODO

sycl/include/CL/sycl/handler.hpp

kbobrovs · 2020-11-12T09:46:18Z

@dm-vodopyanov, I removed the ABI breaking change - please take a look.
I had to rebase atop other patches to have a chance to pass testing. Please review
[SYCL][ESIMD] RT: remove wrapping buffer objects into images. … eca13ef
and later

kbobrovs · 2020-11-14T08:50:53Z

Collapsed all commits into the 4 - 2xFE, RT and the LowerESIMD. That's how they should be committed (w/o further squashing)
I think I addressed all review comments.
@Fznamznon - please check/approve the FE (6fb2ce1, 06dd08c)
@DenisBakhvalov - please check/approve the LowerESIMD (873d4a6)
@dm-vodopyanov - please check/approve the RT (055f859)

Fznamznon

Is there any chance to merge patches for different components separately? It would simplify "git blaming" in the future.

clang/include/clang/Basic/AttrDocs.td

Fznamznon · 2020-11-16T13:18:18Z

clang/test/CodeGenSYCL/esimd-accessor-ptr-md.cpp

+  cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write> accessorA;
+  cl::sycl::accessor<int, 1, cl::sycl::access::mode::read> accessorB;
+
+  cl::sycl::kernel_single_task<class esimd_kernel>(


Could you please use queue::submit and handler::single_task from the mock header instead? And include the mock header as system header using -internal-isystem . Like here https://github.com/intel/llvm/blob/sycl/clang/test/CodeGenSYCL/stall_enable.cpp

We need to follow single style in FE test and this style must be as close as possible to the way how we use real sycl headers.

I don't see single style in FE tests, so I picked one. But OK, I will try to rework as you suggest.

I don't see single style in FE tests, so I picked on. But OK, I will try to rework as you suggest.

Yes, we don't have single style yet. This problem is already captured, and I remember that someone was working on it, but it is not finished.

clang/test/CodeGenSYCL/esimd-accessor-ptr-md.cpp

Fznamznon · 2020-11-16T13:22:14Z

clang/test/CodeGenSYCL/esimd-accessor-ptr-md.cpp

+
+  cl::sycl::kernel_single_task<class esimd_kernel>(
+      [=]() __attribute__((sycl_explicit_simd)) {
+        accessorA.use(val);


Could you please also try other kinds of arguments (like POD for example) and check that all metadata is correct.

Is this a nit? The main point of the test is checking accessors argument translation and MD generation. PODs are no different from other types in this regard - and int val should be enough.

My point was to check that we don't generate something unnecessary for non-accessor types in case if it can confuse underlying backend. And I probably missed that you already have captured int, so yes, this comment is more a nit.

Fznamznon · 2020-11-16T13:23:06Z

clang/test/CodeGenSYCL/esimd-accessor-ptr-md.cpp

+  // CHECK:   i32 "VCArgumentDesc" "VCArgumentIOKind"="0" "VCArgumentKind"="0" %_arg_1,
+  // CHECK:   i32 addrspace(1)* nocapture "VCArgumentDesc"="buffer_t" "VCArgumentIOKind"="0" "VCArgumentKind"="2" %_arg_3)
+  // --- Attributes
+  // CHECK: {{.*}} !kernel_arg_accessor_ptr ![[ACC_PTR_ATTR:[0-9]+]] !sycl_explicit_simd !{{[0-9]+}} {{.*}}{


Could you please check that __init_esimd is called instead of regular __init in case of ESIMD kernel?

This is basically proven by the fact that an accessor is translated into a single pointer arg (as in __init_esimd) rather than 4 (as in __init). I'll see if I can make direct check. Do you have suggestions, BTW?

I don't remember exactly, but I was under impression that we don't enable early optimizations for ESIMD mode, am I right?
In case if IR is not optimized, I would just check that there are two calls to __init_esimd function, i.e. (the following LLVM IR may contain errors):

CHECK: call spir_func void @{{.*}}__init_esimd{{.*}}(%"class{{.*}}cl::sycl::accessor" addrspace(4)* {{.*}}, i32 addrspace(1)* {{.*}}) /* two times */

Or you can even check which argument of the kernel is passed like here

llvm/clang/test/CodeGenSYCL/accessor_inheritance.cpp

Line 84 in 3b8af22

// CHECK: [[GEP2:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class._ZTSZ4mainE3$_0.anon", %"class._ZTSZ4mainE3$_0.anon"* [[KERNEL_OBJ]], i32 0, i32 0

.

In case if IR is optimized for ESIMD mode, we can check presence of __init_esimd in non-optimized IR.

ok, thanks. Will try

Fznamznon · 2020-11-16T13:25:38Z

clang/test/CodeGenSYCL/int_header_esimd.cpp

+// -- ESIMD Lambda kernel.
+template <typename ID, typename F>
+void kernelA(F f) __attribute__((sycl_kernel)) {
+  f();
+}


Same here, please use queue::submit and single_task from the mock header included like system one.

kbobrovs · 2020-11-16T13:45:22Z

Is there any chance to merge patches for different components separately? It would simplify "git blaming" in the future.

that's the intent. I would like this PR to be committed w/o squashing - as 4 patches.

This attribute is to be used by the FE to mark ESIMD kernel arguments originating from buffer accessors, which is then translated to BE-specific metadata needed for correct code generation. Signed-off-by: Konstantin S Bobrovsky <[email protected]>

- Use 'esimd_acc_ptr' attribute to mark ESIMD kernel pointer arguments originating from accessors - Use '__init_esimd' initializer in FE for ESIMD kernel accessors. - Add 'isESIMD()' function to the integration header for the RT to be able to distinguish ESIMD kernels from normal ones for the purpose of proper accessor arg setting. Signed-off-by: Konstantin S Bobrovsky <[email protected]>

kbobrovs · 2020-11-16T16:40:47Z

@Fznamznon - comments are addressed, please take a look. Had to force-push to maintain the proper patch structure, hope this is OK.
@dm-vodopyanov, please re-approve the RT part.

kbobrovs · 2020-11-16T19:40:48Z

@DenisBakhvalov, please review/approve the LowerESIMD part

kbobrovs · 2020-11-17T09:57:09Z

@dm-vodopyanov - could you please restore your approval?
@bader, can you please take a look and merge if no objections?

bader · 2020-11-17T11:09:57Z

@dm-vodopyanov - could you please restore your approval?
@bader, can you please take a look and merge if no objections?

@kbobrovs, I see a few other PR with similar changes. Could you resolve duplication by closing PRs, which are not intended for merging, please? Make sure that all comments are addressed.
In particular: #2699 and #2748.

I have a couple of comments WRT LLVM pass changes (mentioned in #2748):

There is no regression test for LLVM pass
I don't quite get the protocol between FE and LLVM pass and I hoped that we can simplify the logic in the pass. Could you answer this question, please?

kbobrovs · 2020-11-17T12:19:49Z

@bader - I closed #2747 and #2748. #2699 is not part of this one and still needs to be worked on/merged separately.

There is no regression test for LLVM pass

There actually is, but it is part of SYCL CodeGen: clang/test/CodeGenSYCL/esimd-accessor-ptr-md.cpp.
Will add one more to llvm/test/SYCLLowerIR

I don't quite get the protocol between FE and LLVM pass and I hoped that we can simplify the logic in the pass. Could you answer this question, please?

sorry, missed that one. The question is:

Is it possible to have a kernel_arg_accessor_ptr metadata with 0 value?

yes. Only pointer arguments originating from accessors have 1 in the corresponding place in kernel_arg_accessor_ptr metadata

Make sure that all comments are addressed.

I think all are addressed - reviewers all approved. Do you have some specific comment in mind (except the above one)?

kbobrovs · 2020-11-17T14:38:54Z

@bader - test is added, please take a look

bader

LGTM, except one nit.

bader · 2020-11-17T14:44:36Z

llvm/test/SYCLLowerIR/esimd_acc_ptr.ll

+; 'kernel_arg_accessor_ptr' metadata. Particularly, that it generates additional
+; vector of per-argument metadata (accessible from "genx.kernels" top-level
+; metadata node):
+; - for those arguments having '1' in the corresponding


The code seems to allow any non-zero value. Please, either update the comment or add an assert (or some other check) that FE emitted expected value.

implemented the nit

…ents. Signed-off-by: Konstantin S Bobrovsky <[email protected]>

BE is now able to accept buffers instead of image1d buffers, so remove the w/a wrapping. This patch also gets rid of majority of __SYCL_EXPLICIT_SIMD__ macro checks in the API code, paving a way for ESIMD/SYCL kernel co-existence. Details: - Use new int header feature KernelInfo<>::isESIMD() to properly set accessor args - Add __init_esimd accessor methods to be used by FE with ESIMD kernels - Some ESIMD-related stuff is now unsed - can't be removed now not to break ABI. TODOs added. Signed-off-by: Konstantin S Bobrovsky <[email protected]>

kbobrovs · 2020-11-17T15:54:42Z

@bader, please do not squash when merging

bader · 2020-11-17T15:58:44Z

@bader, please do not squash when merging

All other merge options are disabled for this project.
The only way to avoid squash is to create a separate PR for each commit.

kbobrovs · 2020-11-17T16:18:02Z

All other merge options are disabled for this project.
The only way to avoid squash is to create a separate PR for each commit.

That's unfortunate. Those PRs won't pass testing. Can we change this setting BTW?
Otherwise, let's squash then, as we don't seem to have a choice. The description can be used as the commit message.

/SPIRV-LLVM-Translator/lib/SPIRV/SPIRVRegularizeLLVM.cpp:535:15: warning: implicit conversion turns string literal into bool: 'const char[39]' to 'bool' [-Wstring-conversion] 535 | assert(!"Cache controls must decorate a pointer"); | ~^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ /usr/include/assert.h:93:27: note: expanded from macro 'assert' 93 | (static_cast <bool> (expr) \ | ^~~~ The original code is not wrong but `false &&` is only a few characters more and does the same thing without warning. Original commit: KhronosGroup/SPIRV-LLVM-Translator@573e951a3207fe9

kbobrovs requested a review from a team as a code owner November 9, 2020 15:09

kbobrovs requested a review from dm-vodopyanov November 9, 2020 15:09

This was referenced Nov 9, 2020

[SYCL][ESIMD] LowerESIMD: add 'buffer_t' MD for accessor kernel args. #2748

Closed

[SYCL][ESIMD] Refactor ESIMD kernel support in FE. #2747

Closed

dm-vodopyanov reviewed Nov 9, 2020

View reviewed changes

sycl/include/CL/sycl/handler.hpp Outdated Show resolved Hide resolved

DenisBakhvalov reviewed Nov 10, 2020

View reviewed changes

kbobrovs force-pushed the drop_image_wrap_rt branch from b19d7a8 to a689365 Compare November 12, 2020 09:30

kbobrovs requested review from bader, elizabethandrews, Fznamznon and premanandrao as code owners November 12, 2020 09:30

kbobrovs requested a review from dm-vodopyanov November 12, 2020 09:46

kbobrovs force-pushed the drop_image_wrap_rt branch 4 times, most recently from 021fe9a to 055f859 Compare November 14, 2020 08:31

kbobrovs changed the title ~~[SYCL][ESIMD] RT: remove wrapping buffer objects into images.~~ [SYCL][ESIMD] Remove wrapping buffer objects into images. Nov 14, 2020

kbobrovs requested a review from DenisBakhvalov November 16, 2020 07:01

dm-vodopyanov previously approved these changes Nov 16, 2020

View reviewed changes

Fznamznon reviewed Nov 16, 2020

View reviewed changes

kbobrovs dismissed dm-vodopyanov’s stale review via 35c05fa November 16, 2020 13:47

kbobrovs force-pushed the drop_image_wrap_rt branch from 35c05fa to 012a0ef Compare November 16, 2020 16:30

kbobrovs force-pushed the drop_image_wrap_rt branch from 012a0ef to 61d87bb Compare November 16, 2020 16:37

Fznamznon previously approved these changes Nov 16, 2020

View reviewed changes

kbobrovs requested a review from dm-vodopyanov November 16, 2020 19:39

DenisBakhvalov previously approved these changes Nov 16, 2020

View reviewed changes

dm-vodopyanov previously approved these changes Nov 17, 2020

View reviewed changes

kbobrovs dismissed stale reviews from dm-vodopyanov, DenisBakhvalov, and Fznamznon via 079fe1e November 17, 2020 14:35

kbobrovs force-pushed the drop_image_wrap_rt branch from 61d87bb to 079fe1e Compare November 17, 2020 14:35

bader reviewed Nov 17, 2020

View reviewed changes

kbobrovs added 2 commits November 17, 2020 07:49

[SYCL][ESIMD] LowerESIMD: add 'buffer_t' MD for accessor kernel argum…

69d8d7a

…ents. Signed-off-by: Konstantin S Bobrovsky <[email protected]>

kbobrovs force-pushed the drop_image_wrap_rt branch from 079fe1e to cb221b6 Compare November 17, 2020 15:50

kbobrovs requested a review from bader November 17, 2020 15:50

bader approved these changes Nov 17, 2020

View reviewed changes

bader merged commit d2d20d6 into intel:sycl Nov 17, 2020

kbobrovs deleted the drop_image_wrap_rt branch November 17, 2020 17:02

[SYCL][ESIMD] Remove wrapping buffer objects into images. #2746

[SYCL][ESIMD] Remove wrapping buffer objects into images. #2746

Uh oh!

Conversation

kbobrovs commented Nov 9, 2020 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

dm-vodopyanov left a comment • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

kbobrovs commented Nov 12, 2020

Uh oh!

kbobrovs commented Nov 14, 2020 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

Fznamznon left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

kbobrovs Nov 16, 2020 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Fznamznon Nov 16, 2020 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

kbobrovs commented Nov 16, 2020

Uh oh!

kbobrovs commented Nov 16, 2020

Uh oh!

kbobrovs commented Nov 16, 2020

Uh oh!

kbobrovs commented Nov 17, 2020

Uh oh!

bader commented Nov 17, 2020

Uh oh!

kbobrovs commented Nov 17, 2020

Uh oh!

kbobrovs commented Nov 17, 2020

Uh oh!

bader left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

kbobrovs commented Nov 17, 2020

Uh oh!

kbobrovs commented Nov 9, 2020 •

edited

Loading

dm-vodopyanov left a comment •

edited

Loading

kbobrovs commented Nov 14, 2020 •

edited

Loading

kbobrovs Nov 16, 2020 •

edited

Loading

Fznamznon Nov 16, 2020 •

edited

Loading