[SYCL][ESIMD] Use vload/vstore for simd object getter/setter. #1887

kbobrovs · 2020-06-13T02:07:25Z

DO NOT REVIEW FIRST TWO COMMITS (covered by #1881)

Use special intrinsics to access simd objects in the private address space
to disable standard LLVM optimizations on them.

Author: Gang Chen [email protected]

Signed-off-by: Konstantin S Bobrovsky [email protected]

kbobrovs · 2020-06-13T02:08:20Z

@cmc-rep, please pick the PR as the author of the code

againull · 2020-06-17T18:12:25Z

llvm/test/SYCLLowerIR/esimd_lower_ldst.ll

+; Function Attrs: norecurse nounwind
+define dso_local spir_func void @_Z3foov() local_unnamed_addr #1 {
+; CHECK-LABEL: @_Z3foov(
+; CHECK-NEXT: [[TMP1:%.*]] = call <16 x i32> @llvm.genx.vload.v16i32.p0v16i32(<16 x i32>* getelementptr inbounds (%"class._ZTSN4sycl5intel3gpu4simdIiLi16EEE.sycl::intel::gpu::simd", %"class._ZTSN4sycl5intel3gpu4simdIiLi16EEE.sycl::intel::gpu::simd"* @vg, i64 0, i32 0))


Is it expected that vload is not lowered in this case? It is not obvious why only vstore is lowered to the normal store.

Yes, this is we are testing. vload/vstore associated with global variable with special attritube stays as vload/vstore. vload/vstore not associated with those global variables is changed to generic load/store. This is explained in the code comment

This means that the lowering of the load is not tested, please add additional test to check lowering of the load.

load/store are processed in the same fashion. I added one case for load and one case for store. That should be sufficient

I don't see a test for load lowering, how I can be sure that your code for load lowering is working?

againull · 2020-06-17T18:15:14Z

llvm/lib/SYCLLowerIR/LowerESIMDVLoadVStore.cpp

+
+bool ESIMDLowerLoadStore::runOnFunction(Function &F) {
+  bool Changed = false;
+  Changed |= lowerLoadStore(F);


Is it possible to replace these 3 lines with
return lowerLoadStore(F); ?

Why. It works and it is correct. Do we really need to pick on this detail?

Is this review only about functional correctness?

That is fine. I can change this one.

againull · 2020-06-17T18:17:04Z

llvm/lib/SYCLLowerIR/LowerESIMDVLoadVStore.cpp

+#include "llvm/IR/Module.h"
+#include "llvm/InitializePasses.h"
+#include "llvm/Support/Debug.h"
+#include "llvm/Support/raw_ostream.h"


Probably some includes are not needed.

Which includes are not needed. Please be specific

Review process doesn't work this way.
I pointed to "#include "llvm/Support/raw_ostream.h" which probably was used during debugging the code but is not needed in the published code. And you should check if other headers are needed.

I will delete raw_ostream.h

againull · 2020-06-17T18:29:27Z

llvm/include/llvm/SYCLLowerIR/LowerESIMD.h

@@ -34,6 +34,9 @@ class SYCLLowerESIMDPass : public PassInfoMixin<SYCLLowerESIMDPass> {
 FunctionPass *createSYCLLowerESIMDPass();
 void initializeSYCLLowerESIMDLegacyPassPass(PassRegistry &);

+FunctionPass *createESIMDLowerLoadStorePass();


Pass is written to work with old pass manager only. Please support it for the new pass manager, SYCLLowerESIMDPass or other pass can be used as an example.

againull · 2020-06-17T18:31:49Z

llvm/lib/SYCLLowerIR/LowerESIMDVLoadVStore.cpp

+  // lower all vload/vstore into normal load/store.
+  std::vector<Instruction *> ToErase;
+  for (Instruction &Inst : instructions(F)) {
+    if (GenXIntrinsic::isVLoadStore(&Inst)) {


This loop body is one huge block which is hard to read.

Could you please use more early exits there? I think it will be much easier to read. For example:

if (!GenXIntrinsic::isVLoadStore(&Inst)) continue;

and so on.

I think this is just your personal preference. Other people may prefer not to use continue

This is not my personal preference. Here is the link to the LLVM coding guidelines:
https://llvm.org/docs/CodingStandards.html#use-early-exits-and-continue-to-simplify-code

Coding style is just a guideline. Individual developer should have some freedom in applying it. In this case, for example, I do not like the negative sense in the condition when using condition. I prefer to use positive sense whenever possible for the condition test. Clearly, you read this code, and it did not confuse you.

Individual reviewers should have freedom in insisting to apply such guidelines (which are essentially a set of BKMs to write robust/correct/maintainable code) if they seem them suitable. I'm exercising that freedom.

Ok, I understand that you disagree, but I am not the person who determines the review process for this project.
@pvchupin @olegmaslovatintel @bader @romanovvlad Could you please help with this question?

This project follows LLVM coding guidelines and this is requirement.

I have used "continue"

Probably you forgot to upload the new changes.

Not sure what I did wrong. It should be up to date now: https://github.com/kbobrovs/llvm/blob/esimd-ldst/llvm/lib/SYCLLowerIR/LowerESIMDVLoadVStore.cpp

kbobrovs · 2020-07-01T23:12:03Z

@cmc-rep, please check once again if review comments are addressed.
@againull, can you please take another look and see if there are unaddressed comments.
I had to rebase and force-push to resolve conflicts.

cmc-rep · 2020-07-02T04:53:50Z

I have addressed all the comments last week

kbobrovs · 2020-07-03T01:00:00Z

I have addressed all the comments last week

@againull, please take another look

Use special intrinsics to access simd objects in the private address space to disable standard LLVM optimizations on them. Author: Gang Chen <[email protected]> Signed-off-by: Konstantin S Bobrovsky <[email protected]>

Signed-off-by: gangche1 <[email protected]>

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

kbobrovs · 2020-07-05T00:15:11Z

Had to rebase after ESIMD lowering merge. @againull, can you please approve again.

List of changes: All Literal parameters of instructions in OpenCL.DebugInfo.100 are OpConstants in NonSemantic.Shader.DebugInfo.100 and NonSemantic.Shader.DebugInfo.200; SPV_KHR_non_semantic_info is being implicitly added for nonsemantic debug info; Original commit: KhronosGroup/SPIRV-LLVM-Translator@3fe4393

kbobrovs added the esimd Explicit SIMD feature label Jun 13, 2020

kbobrovs requested a review from bader as a code owner June 13, 2020 02:07

kbobrovs requested review from againull and AlexeySachkov June 13, 2020 02:07

againull requested changes Jun 17, 2020

View reviewed changes

cmc-rep force-pushed the esimd-ldst branch from 6190e66 to 61eae7c Compare June 24, 2020 16:03

kbobrovs force-pushed the esimd-ldst branch from 61eae7c to 3146cdd Compare July 1, 2020 23:09

kbobrovs requested a review from a team as a code owner July 1, 2020 23:09

againull previously approved these changes Jul 3, 2020

View reviewed changes

kbobrovs and others added 3 commits July 4, 2020 17:09

[SYCL][ESIMD] Use vload/vstore for simd object getter/setter.

8f54872

Use special intrinsics to access simd objects in the private address space to disable standard LLVM optimizations on them. Author: Gang Chen <[email protected]> Signed-off-by: Konstantin S Bobrovsky <[email protected]>

[SYCL][ESIMD] - revise vload-vstore lowering

e4d3764

Signed-off-by: gangche1 <[email protected]>

[SQUASH] Apply clang-format.

e7ec678

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

kbobrovs dismissed againull’s stale review via e7ec678 July 5, 2020 00:13

kbobrovs force-pushed the esimd-ldst branch from a51b284 to e7ec678 Compare July 5, 2020 00:13

bader requested a review from againull July 5, 2020 14:44

againull approved these changes Jul 6, 2020

View reviewed changes

bader approved these changes Jul 8, 2020

View reviewed changes

bader merged commit ee98d31 into intel:sycl Jul 8, 2020

kbobrovs deleted the esimd-ldst branch July 30, 2020 12:30

[SYCL][ESIMD] Use vload/vstore for simd object getter/setter. #1887

[SYCL][ESIMD] Use vload/vstore for simd object getter/setter. #1887

Uh oh!

Conversation

kbobrovs commented Jun 13, 2020

Uh oh!

kbobrovs commented Jun 13, 2020

Uh oh!

againull Jun 17, 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!

againull Jun 17, 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!

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!

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!

Choose a reason for hiding this comment

Uh oh!

againull Jun 22, 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!

Choose a reason for hiding this comment

Uh oh!

kbobrovs commented Jul 1, 2020

Uh oh!

cmc-rep commented Jul 2, 2020

Uh oh!

kbobrovs commented Jul 3, 2020

Uh oh!

kbobrovs commented Jul 5, 2020

Uh oh!

Uh oh!

againull Jun 17, 2020 •

edited

Loading

againull Jun 17, 2020 •

edited

Loading

againull Jun 22, 2020 •

edited

Loading